diff --git a/src/.gitignore b/src/.gitignore index 13fb446225e0c84c6d51e6e1b37da9303fd10b3f..fe23bc1f55446fc6873390ef1fca32af96d94d8d 100644 --- a/src/.gitignore +++ b/src/.gitignore @@ -25,6 +25,7 @@ /kokkos.h /kokkos_type.h /kokkos_few.h +/kokkos_base.h /manifold*.cpp /manifold*.h diff --git a/src/GPU/pair_eam_alloy_gpu.cpp b/src/GPU/pair_eam_alloy_gpu.cpp index ab0f499a98e6813a875753d63302cfe34bd8fa67..9b42b0a14fab56b44652942973fd122a7fe113da 100644 --- a/src/GPU/pair_eam_alloy_gpu.cpp +++ b/src/GPU/pair_eam_alloy_gpu.cpp @@ -28,6 +28,7 @@ #include "error.h" #include "neigh_request.h" #include "gpu_extra.h" +#include "domain.h" using namespace LAMMPS_NS; diff --git a/src/GPU/pair_eam_fs_gpu.cpp b/src/GPU/pair_eam_fs_gpu.cpp index a2b339db9ac6ce197b19ad7705edfb24ea132599..c29b49631ce210121df7479b3fc60867bca31438 100644 --- a/src/GPU/pair_eam_fs_gpu.cpp +++ b/src/GPU/pair_eam_fs_gpu.cpp @@ -28,6 +28,7 @@ #include "error.h" #include "neigh_request.h" #include "gpu_extra.h" +#include "domain.h" using namespace LAMMPS_NS; diff --git a/src/GRANULAR/fix_wall_gran.cpp b/src/GRANULAR/fix_wall_gran.cpp index eeec94fdf729a76694188a715119781adfaaefe5..033c35dbacc0e96bb910cd1c3e6b2fa261eb815a 100644 --- a/src/GRANULAR/fix_wall_gran.cpp +++ b/src/GRANULAR/fix_wall_gran.cpp @@ -30,6 +30,7 @@ #include "math_const.h" #include "memory.h" #include "error.h" +#include "neighbor.h" using namespace LAMMPS_NS; using namespace FixConst; diff --git a/src/GRANULAR/fix_wall_gran_region.cpp b/src/GRANULAR/fix_wall_gran_region.cpp index a09b9dfa421b971df17992722ccc55c1308b1454..d1c5d4c9c771f8a73bd0d95165ee362976f50509 100644 --- a/src/GRANULAR/fix_wall_gran_region.cpp +++ b/src/GRANULAR/fix_wall_gran_region.cpp @@ -30,6 +30,8 @@ #include "math_const.h" #include "memory.h" #include "error.h" +#include "comm.h" +#include "neighbor.h" using namespace LAMMPS_NS; using namespace FixConst; diff --git a/src/KOKKOS/Install.sh b/src/KOKKOS/Install.sh index a382295be3556580921b86b6184c7fa1a6cd035a..295e46bcaec54282a9511adb4ff5af91a9b75fff 100644 --- a/src/KOKKOS/Install.sh +++ b/src/KOKKOS/Install.sh @@ -28,8 +28,20 @@ action () { # force rebuild of files with LMP_KOKKOS switch -touch ../accelerator_kokkos.h -touch ../memory.h +KOKKOS_INSTALLED=0 +if (test -e ../Makefile.package) then + KOKKOS_INSTALLED=`grep DLMP_KOKKOS ../Makefile.package | wc -l` +fi + +if (test $mode = 1) then + if (test $KOKKOS_INSTALLED = 0) then + touch ../accelerator_kokkos.h + fi +elif (test $mode = 0) then + if (test $KOKKOS_INSTALLED = 1) then + touch ../accelerator_kokkos.h + fi +fi # list of files with optional dependcies @@ -125,8 +137,9 @@ action improper_harmonic_kokkos.cpp improper_harmonic.cpp action improper_harmonic_kokkos.h improper_harmonic.h action kokkos.cpp action kokkos.h -action kokkos_type.h +action kokkos_base.h action kokkos_few.h +action kokkos_type.h action memory_kokkos.h action modify_kokkos.cpp action modify_kokkos.h diff --git a/src/KOKKOS/angle_charmm_kokkos.cpp b/src/KOKKOS/angle_charmm_kokkos.cpp index 401a00c856ffcc86543f22024bbb0cec1d9dc2e3..59a20c25df36184d4a23a9bf9f6071c625f5eb38 100644 --- a/src/KOKKOS/angle_charmm_kokkos.cpp +++ b/src/KOKKOS/angle_charmm_kokkos.cpp @@ -24,7 +24,7 @@ #include "comm.h" #include "force.h" #include "math_const.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -51,8 +51,8 @@ template<class DeviceType> AngleCharmmKokkos<DeviceType>::~AngleCharmmKokkos() { if (!copymode) { - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); } } @@ -71,15 +71,15 @@ void AngleCharmmKokkos<DeviceType>::compute(int eflag_in, int vflag_in) if (eflag_atom) { //if(k_eatom.dimension_0()<maxeatom) { // won't work without adding zero functor - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom"); d_eatom = k_eatom.template view<DeviceType>(); //} } if (vflag_atom) { //if(k_vatom.dimension_0()<maxvatom) { // won't work without adding zero functor - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"improper:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"improper:vatom"); d_vatom = k_vatom.template view<DeviceType>(); //} } diff --git a/src/KOKKOS/angle_class2_kokkos.cpp b/src/KOKKOS/angle_class2_kokkos.cpp index e851e9d500b10246571210a6c38daae5d472528a..108b4f48f225827228fe1a5667edbd7a7c8ad6aa 100644 --- a/src/KOKKOS/angle_class2_kokkos.cpp +++ b/src/KOKKOS/angle_class2_kokkos.cpp @@ -24,7 +24,7 @@ #include "comm.h" #include "force.h" #include "math_const.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -51,8 +51,8 @@ template<class DeviceType> AngleClass2Kokkos<DeviceType>::~AngleClass2Kokkos() { if (!copymode) { - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); } } @@ -70,13 +70,13 @@ void AngleClass2Kokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"angle:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"angle:eatom"); d_eatom = k_eatom.template view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"angle:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"angle:vatom"); d_vatom = k_vatom.template view<DeviceType>(); } diff --git a/src/KOKKOS/angle_harmonic_kokkos.cpp b/src/KOKKOS/angle_harmonic_kokkos.cpp index 9fd237ddb3e90404aa9e9a979be2ccaba3953e6d..dd5a1e26c7c55d135a5e37a0619dace64b49dbb1 100644 --- a/src/KOKKOS/angle_harmonic_kokkos.cpp +++ b/src/KOKKOS/angle_harmonic_kokkos.cpp @@ -24,7 +24,7 @@ #include "comm.h" #include "force.h" #include "math_const.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -51,8 +51,8 @@ template<class DeviceType> AngleHarmonicKokkos<DeviceType>::~AngleHarmonicKokkos() { if (!copymode) { - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); } } @@ -70,13 +70,13 @@ void AngleHarmonicKokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"angle:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"angle:eatom"); d_eatom = k_eatom.template view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"angle:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"angle:vatom"); d_vatom = k_vatom.template view<DeviceType>(); } diff --git a/src/KOKKOS/atom_kokkos.cpp b/src/KOKKOS/atom_kokkos.cpp index 31b33dbdc9a2c406e40e6e36f8898d88d1c50d5b..4ecead5b1d052be975c40069d169a7fad43cbf63 100644 --- a/src/KOKKOS/atom_kokkos.cpp +++ b/src/KOKKOS/atom_kokkos.cpp @@ -19,7 +19,7 @@ #include "update.h" #include "domain.h" #include "atom_masks.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "kokkos.h" @@ -33,59 +33,59 @@ AtomKokkos::AtomKokkos(LAMMPS *lmp) : Atom(lmp) {} AtomKokkos::~AtomKokkos() { - memory->destroy_kokkos(k_tag, tag); - memory->destroy_kokkos(k_mask, mask); - memory->destroy_kokkos(k_type, type); - memory->destroy_kokkos(k_image, image); - memory->destroy_kokkos(k_molecule, molecule); - - memory->destroy_kokkos(k_x, x); - memory->destroy_kokkos(k_v, v); - memory->destroy_kokkos(k_f, f); - - memory->destroy_kokkos(k_mass, mass); - memory->destroy_kokkos(k_q, q); - - memory->destroy_kokkos(k_radius, radius); - memory->destroy_kokkos(k_rmass, rmass); - memory->destroy_kokkos(k_omega, omega); - memory->destroy_kokkos(k_angmom, angmom); - memory->destroy_kokkos(k_torque, torque); - - memory->destroy_kokkos(k_nspecial, nspecial); - memory->destroy_kokkos(k_special, special); - memory->destroy_kokkos(k_num_bond, num_bond); - memory->destroy_kokkos(k_bond_type, bond_type); - memory->destroy_kokkos(k_bond_atom, bond_atom); - memory->destroy_kokkos(k_num_angle, num_angle); - memory->destroy_kokkos(k_angle_type, angle_type); - memory->destroy_kokkos(k_angle_atom1, angle_atom1); - memory->destroy_kokkos(k_angle_atom2, angle_atom2); - memory->destroy_kokkos(k_angle_atom3, angle_atom3); - memory->destroy_kokkos(k_num_dihedral, num_dihedral); - memory->destroy_kokkos(k_dihedral_type, dihedral_type); - memory->destroy_kokkos(k_dihedral_atom1, dihedral_atom1); - memory->destroy_kokkos(k_dihedral_atom2, dihedral_atom2); - memory->destroy_kokkos(k_dihedral_atom3, dihedral_atom3); - memory->destroy_kokkos(k_dihedral_atom4, dihedral_atom4); - memory->destroy_kokkos(k_num_improper, num_improper); - memory->destroy_kokkos(k_improper_type, improper_type); - memory->destroy_kokkos(k_improper_atom1, improper_atom1); - memory->destroy_kokkos(k_improper_atom2, improper_atom2); - memory->destroy_kokkos(k_improper_atom3, improper_atom3); - memory->destroy_kokkos(k_improper_atom4, improper_atom4); + memoryKK->destroy_kokkos(k_tag, tag); + memoryKK->destroy_kokkos(k_mask, mask); + memoryKK->destroy_kokkos(k_type, type); + memoryKK->destroy_kokkos(k_image, image); + memoryKK->destroy_kokkos(k_molecule, molecule); + + memoryKK->destroy_kokkos(k_x, x); + memoryKK->destroy_kokkos(k_v, v); + memoryKK->destroy_kokkos(k_f, f); + + memoryKK->destroy_kokkos(k_mass, mass); + memoryKK->destroy_kokkos(k_q, q); + + memoryKK->destroy_kokkos(k_radius, radius); + memoryKK->destroy_kokkos(k_rmass, rmass); + memoryKK->destroy_kokkos(k_omega, omega); + memoryKK->destroy_kokkos(k_angmom, angmom); + memoryKK->destroy_kokkos(k_torque, torque); + + memoryKK->destroy_kokkos(k_nspecial, nspecial); + memoryKK->destroy_kokkos(k_special, special); + memoryKK->destroy_kokkos(k_num_bond, num_bond); + memoryKK->destroy_kokkos(k_bond_type, bond_type); + memoryKK->destroy_kokkos(k_bond_atom, bond_atom); + memoryKK->destroy_kokkos(k_num_angle, num_angle); + memoryKK->destroy_kokkos(k_angle_type, angle_type); + memoryKK->destroy_kokkos(k_angle_atom1, angle_atom1); + memoryKK->destroy_kokkos(k_angle_atom2, angle_atom2); + memoryKK->destroy_kokkos(k_angle_atom3, angle_atom3); + memoryKK->destroy_kokkos(k_num_dihedral, num_dihedral); + memoryKK->destroy_kokkos(k_dihedral_type, dihedral_type); + memoryKK->destroy_kokkos(k_dihedral_atom1, dihedral_atom1); + memoryKK->destroy_kokkos(k_dihedral_atom2, dihedral_atom2); + memoryKK->destroy_kokkos(k_dihedral_atom3, dihedral_atom3); + memoryKK->destroy_kokkos(k_dihedral_atom4, dihedral_atom4); + memoryKK->destroy_kokkos(k_num_improper, num_improper); + memoryKK->destroy_kokkos(k_improper_type, improper_type); + memoryKK->destroy_kokkos(k_improper_atom1, improper_atom1); + memoryKK->destroy_kokkos(k_improper_atom2, improper_atom2); + memoryKK->destroy_kokkos(k_improper_atom3, improper_atom3); + memoryKK->destroy_kokkos(k_improper_atom4, improper_atom4); // USER-DPD package - memory->destroy_kokkos(k_uCond,uCond); - memory->destroy_kokkos(k_uMech,uMech); - memory->destroy_kokkos(k_uChem,uChem); - memory->destroy_kokkos(k_uCG,uCG); - memory->destroy_kokkos(k_uCGnew,uCGnew); - memory->destroy_kokkos(k_rho,rho); - memory->destroy_kokkos(k_dpdTheta,dpdTheta); - memory->destroy_kokkos(k_duChem,duChem); - - memory->destroy_kokkos(k_dvector,dvector); + memoryKK->destroy_kokkos(k_uCond,uCond); + memoryKK->destroy_kokkos(k_uMech,uMech); + memoryKK->destroy_kokkos(k_uChem,uChem); + memoryKK->destroy_kokkos(k_uCG,uCG); + memoryKK->destroy_kokkos(k_uCGnew,uCGnew); + memoryKK->destroy_kokkos(k_rho,rho); + memoryKK->destroy_kokkos(k_dpdTheta,dpdTheta); + memoryKK->destroy_kokkos(k_duChem,duChem); + + memoryKK->destroy_kokkos(k_dvector,dvector); dvector = NULL; } @@ -232,10 +232,10 @@ void AtomKokkos::sort() void AtomKokkos::grow(unsigned int mask){ if (mask & SPECIAL_MASK){ - memory->destroy_kokkos(k_special, special); + memoryKK->destroy_kokkos(k_special, special); sync(Device, mask); modified(Device, mask); - memory->grow_kokkos(k_special,special,nmax,maxspecial,"atom:special"); + memoryKK->grow_kokkos(k_special,special,nmax,maxspecial,"atom:special"); avec->grow_reset(); sync(Host, mask); } @@ -270,7 +270,7 @@ int AtomKokkos::add_custom(const char *name, int flag) int n = strlen(name) + 1; dname[index] = new char[n]; strcpy(dname[index],name); - memory->grow_kokkos(k_dvector,dvector,ndvector,nmax, + memoryKK->grow_kokkos(k_dvector,dvector,ndvector,nmax, "atom:dvector"); } @@ -291,7 +291,7 @@ void AtomKokkos::remove_custom(int flag, int index) delete [] iname[index]; iname[index] = NULL; } else { - //memory->destroy_kokkos(dvector); + //memoryKK->destroy_kokkos(dvector); dvector[index] = NULL; delete [] dname[index]; dname[index] = NULL; @@ -302,25 +302,25 @@ void AtomKokkos::remove_custom(int flag, int index) void AtomKokkos::deallocate_topology() { - memory->destroy_kokkos(k_bond_type, bond_type); - memory->destroy_kokkos(k_bond_atom, bond_atom); - - memory->destroy_kokkos(k_angle_type, angle_type); - memory->destroy_kokkos(k_angle_atom1, angle_atom1); - memory->destroy_kokkos(k_angle_atom2, angle_atom2); - memory->destroy_kokkos(k_angle_atom3, angle_atom3); - - memory->destroy_kokkos(k_dihedral_type, dihedral_type); - memory->destroy_kokkos(k_dihedral_atom1, dihedral_atom1); - memory->destroy_kokkos(k_dihedral_atom2, dihedral_atom2); - memory->destroy_kokkos(k_dihedral_atom3, dihedral_atom3); - memory->destroy_kokkos(k_dihedral_atom4, dihedral_atom4); - - memory->destroy_kokkos(k_improper_type, improper_type); - memory->destroy_kokkos(k_improper_atom1, improper_atom1); - memory->destroy_kokkos(k_improper_atom2, improper_atom2); - memory->destroy_kokkos(k_improper_atom3, improper_atom3); - memory->destroy_kokkos(k_improper_atom4, improper_atom4); + memoryKK->destroy_kokkos(k_bond_type, bond_type); + memoryKK->destroy_kokkos(k_bond_atom, bond_atom); + + memoryKK->destroy_kokkos(k_angle_type, angle_type); + memoryKK->destroy_kokkos(k_angle_atom1, angle_atom1); + memoryKK->destroy_kokkos(k_angle_atom2, angle_atom2); + memoryKK->destroy_kokkos(k_angle_atom3, angle_atom3); + + memoryKK->destroy_kokkos(k_dihedral_type, dihedral_type); + memoryKK->destroy_kokkos(k_dihedral_atom1, dihedral_atom1); + memoryKK->destroy_kokkos(k_dihedral_atom2, dihedral_atom2); + memoryKK->destroy_kokkos(k_dihedral_atom3, dihedral_atom3); + memoryKK->destroy_kokkos(k_dihedral_atom4, dihedral_atom4); + + memoryKK->destroy_kokkos(k_improper_type, improper_type); + memoryKK->destroy_kokkos(k_improper_atom1, improper_atom1); + memoryKK->destroy_kokkos(k_improper_atom2, improper_atom2); + memoryKK->destroy_kokkos(k_improper_atom3, improper_atom3); + memoryKK->destroy_kokkos(k_improper_atom4, improper_atom4); } /* ---------------------------------------------------------------------- diff --git a/src/KOKKOS/atom_vec_angle_kokkos.cpp b/src/KOKKOS/atom_vec_angle_kokkos.cpp index 05414cf2e45cdc92f9750c552cb390d915e151cb..a9e55f530a2c1ab3e93eb606f7375bf674c678aa 100644 --- a/src/KOKKOS/atom_vec_angle_kokkos.cpp +++ b/src/KOKKOS/atom_vec_angle_kokkos.cpp @@ -19,7 +19,7 @@ #include "modify.h" #include "fix.h" #include "atom_masks.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" using namespace LAMMPS_NS; @@ -68,33 +68,33 @@ void AtomVecAngleKokkos::grow(int n) sync(Device,ALL_MASK); modified(Device,ALL_MASK); - memory->grow_kokkos(atomKK->k_tag,atomKK->tag,nmax,"atom:tag"); - memory->grow_kokkos(atomKK->k_type,atomKK->type,nmax,"atom:type"); - memory->grow_kokkos(atomKK->k_mask,atomKK->mask,nmax,"atom:mask"); - memory->grow_kokkos(atomKK->k_image,atomKK->image,nmax,"atom:image"); + memoryKK->grow_kokkos(atomKK->k_tag,atomKK->tag,nmax,"atom:tag"); + memoryKK->grow_kokkos(atomKK->k_type,atomKK->type,nmax,"atom:type"); + memoryKK->grow_kokkos(atomKK->k_mask,atomKK->mask,nmax,"atom:mask"); + memoryKK->grow_kokkos(atomKK->k_image,atomKK->image,nmax,"atom:image"); - memory->grow_kokkos(atomKK->k_x,atomKK->x,nmax,3,"atom:x"); - memory->grow_kokkos(atomKK->k_v,atomKK->v,nmax,3,"atom:v"); - memory->grow_kokkos(atomKK->k_f,atomKK->f,nmax,3,"atom:f"); + memoryKK->grow_kokkos(atomKK->k_x,atomKK->x,nmax,3,"atom:x"); + memoryKK->grow_kokkos(atomKK->k_v,atomKK->v,nmax,3,"atom:v"); + memoryKK->grow_kokkos(atomKK->k_f,atomKK->f,nmax,3,"atom:f"); - memory->grow_kokkos(atomKK->k_molecule,atomKK->molecule,nmax,"atom:molecule"); - memory->grow_kokkos(atomKK->k_nspecial,atomKK->nspecial,nmax,3,"atom:nspecial"); - memory->grow_kokkos(atomKK->k_special,atomKK->special,nmax,atomKK->maxspecial, + memoryKK->grow_kokkos(atomKK->k_molecule,atomKK->molecule,nmax,"atom:molecule"); + memoryKK->grow_kokkos(atomKK->k_nspecial,atomKK->nspecial,nmax,3,"atom:nspecial"); + memoryKK->grow_kokkos(atomKK->k_special,atomKK->special,nmax,atomKK->maxspecial, "atom:special"); - memory->grow_kokkos(atomKK->k_num_bond,atomKK->num_bond,nmax,"atom:num_bond"); - memory->grow_kokkos(atomKK->k_bond_type,atomKK->bond_type,nmax,atomKK->bond_per_atom, + memoryKK->grow_kokkos(atomKK->k_num_bond,atomKK->num_bond,nmax,"atom:num_bond"); + memoryKK->grow_kokkos(atomKK->k_bond_type,atomKK->bond_type,nmax,atomKK->bond_per_atom, "atom:bond_type"); - memory->grow_kokkos(atomKK->k_bond_atom,atomKK->bond_atom,nmax,atomKK->bond_per_atom, + memoryKK->grow_kokkos(atomKK->k_bond_atom,atomKK->bond_atom,nmax,atomKK->bond_per_atom, "atom:bond_atom"); - memory->grow_kokkos(atomKK->k_num_angle,atomKK->num_angle,nmax,"atom:num_angle"); - memory->grow_kokkos(atomKK->k_angle_type,atomKK->angle_type,nmax,atomKK->angle_per_atom, + memoryKK->grow_kokkos(atomKK->k_num_angle,atomKK->num_angle,nmax,"atom:num_angle"); + memoryKK->grow_kokkos(atomKK->k_angle_type,atomKK->angle_type,nmax,atomKK->angle_per_atom, "atom:angle_type"); - memory->grow_kokkos(atomKK->k_angle_atom1,atomKK->angle_atom1,nmax,atomKK->angle_per_atom, + memoryKK->grow_kokkos(atomKK->k_angle_atom1,atomKK->angle_atom1,nmax,atomKK->angle_per_atom, "atom:angle_atom1"); - memory->grow_kokkos(atomKK->k_angle_atom2,atomKK->angle_atom2,nmax,atomKK->angle_per_atom, + memoryKK->grow_kokkos(atomKK->k_angle_atom2,atomKK->angle_atom2,nmax,atomKK->angle_per_atom, "atom:angle_atom2"); - memory->grow_kokkos(atomKK->k_angle_atom3,atomKK->angle_atom3,nmax,atomKK->angle_per_atom, + memoryKK->grow_kokkos(atomKK->k_angle_atom3,atomKK->angle_atom3,nmax,atomKK->angle_per_atom, "atom:angle_atom3"); grow_reset(); diff --git a/src/KOKKOS/atom_vec_atomic_kokkos.cpp b/src/KOKKOS/atom_vec_atomic_kokkos.cpp index 6c610c8c111409d8c18201d9be23e44a4fe9195c..f021c45db6feb9c1407115dd1cbeadebf9ddf1d6 100644 --- a/src/KOKKOS/atom_vec_atomic_kokkos.cpp +++ b/src/KOKKOS/atom_vec_atomic_kokkos.cpp @@ -19,7 +19,7 @@ #include "modify.h" #include "fix.h" #include "atom_masks.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" using namespace LAMMPS_NS; @@ -64,14 +64,14 @@ void AtomVecAtomicKokkos::grow(int n) sync(Device,ALL_MASK); modified(Device,ALL_MASK); - memory->grow_kokkos(atomKK->k_tag,atomKK->tag,nmax,"atom:tag"); - memory->grow_kokkos(atomKK->k_type,atomKK->type,nmax,"atom:type"); - memory->grow_kokkos(atomKK->k_mask,atomKK->mask,nmax,"atom:mask"); - memory->grow_kokkos(atomKK->k_image,atomKK->image,nmax,"atom:image"); + memoryKK->grow_kokkos(atomKK->k_tag,atomKK->tag,nmax,"atom:tag"); + memoryKK->grow_kokkos(atomKK->k_type,atomKK->type,nmax,"atom:type"); + memoryKK->grow_kokkos(atomKK->k_mask,atomKK->mask,nmax,"atom:mask"); + memoryKK->grow_kokkos(atomKK->k_image,atomKK->image,nmax,"atom:image"); - memory->grow_kokkos(atomKK->k_x,atomKK->x,nmax,3,"atom:x"); - memory->grow_kokkos(atomKK->k_v,atomKK->v,nmax,3,"atom:v"); - memory->grow_kokkos(atomKK->k_f,atomKK->f,nmax,3,"atom:f"); + memoryKK->grow_kokkos(atomKK->k_x,atomKK->x,nmax,3,"atom:x"); + memoryKK->grow_kokkos(atomKK->k_v,atomKK->v,nmax,3,"atom:v"); + memoryKK->grow_kokkos(atomKK->k_f,atomKK->f,nmax,3,"atom:f"); grow_reset(); sync(Host,ALL_MASK); diff --git a/src/KOKKOS/atom_vec_bond_kokkos.cpp b/src/KOKKOS/atom_vec_bond_kokkos.cpp index 076144420c1442c5db69008905263313b4757be2..bf682c507fd5c3210db20a1e50b35dc09d16a1c7 100644 --- a/src/KOKKOS/atom_vec_bond_kokkos.cpp +++ b/src/KOKKOS/atom_vec_bond_kokkos.cpp @@ -19,7 +19,7 @@ #include "modify.h" #include "fix.h" #include "atom_masks.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" using namespace LAMMPS_NS; @@ -65,21 +65,21 @@ void AtomVecBondKokkos::grow(int n) sync(Device,ALL_MASK); modified(Device,ALL_MASK); - memory->grow_kokkos(atomKK->k_tag,atomKK->tag,nmax,"atom:tag"); - memory->grow_kokkos(atomKK->k_type,atomKK->type,nmax,"atom:type"); - memory->grow_kokkos(atomKK->k_mask,atomKK->mask,nmax,"atom:mask"); - memory->grow_kokkos(atomKK->k_image,atomKK->image,nmax,"atom:image"); - - memory->grow_kokkos(atomKK->k_x,atomKK->x,nmax,3,"atom:x"); - memory->grow_kokkos(atomKK->k_v,atomKK->v,nmax,3,"atom:v"); - memory->grow_kokkos(atomKK->k_f,atomKK->f,nmax,3,"atom:f"); - - memory->grow_kokkos(atomKK->k_molecule,atomKK->molecule,nmax,"atom:molecule"); - memory->grow_kokkos(atomKK->k_nspecial,atomKK->nspecial,nmax,3,"atom:nspecial"); - memory->grow_kokkos(atomKK->k_special,atomKK->special,nmax,atomKK->maxspecial,"atom:special"); - memory->grow_kokkos(atomKK->k_num_bond,atomKK->num_bond,nmax,"atom:num_bond"); - memory->grow_kokkos(atomKK->k_bond_type,atomKK->bond_type,nmax,atomKK->bond_per_atom,"atom:bond_type"); - memory->grow_kokkos(atomKK->k_bond_atom,atomKK->bond_atom,nmax,atomKK->bond_per_atom,"atom:bond_atom"); + memoryKK->grow_kokkos(atomKK->k_tag,atomKK->tag,nmax,"atom:tag"); + memoryKK->grow_kokkos(atomKK->k_type,atomKK->type,nmax,"atom:type"); + memoryKK->grow_kokkos(atomKK->k_mask,atomKK->mask,nmax,"atom:mask"); + memoryKK->grow_kokkos(atomKK->k_image,atomKK->image,nmax,"atom:image"); + + memoryKK->grow_kokkos(atomKK->k_x,atomKK->x,nmax,3,"atom:x"); + memoryKK->grow_kokkos(atomKK->k_v,atomKK->v,nmax,3,"atom:v"); + memoryKK->grow_kokkos(atomKK->k_f,atomKK->f,nmax,3,"atom:f"); + + memoryKK->grow_kokkos(atomKK->k_molecule,atomKK->molecule,nmax,"atom:molecule"); + memoryKK->grow_kokkos(atomKK->k_nspecial,atomKK->nspecial,nmax,3,"atom:nspecial"); + memoryKK->grow_kokkos(atomKK->k_special,atomKK->special,nmax,atomKK->maxspecial,"atom:special"); + memoryKK->grow_kokkos(atomKK->k_num_bond,atomKK->num_bond,nmax,"atom:num_bond"); + memoryKK->grow_kokkos(atomKK->k_bond_type,atomKK->bond_type,nmax,atomKK->bond_per_atom,"atom:bond_type"); + memoryKK->grow_kokkos(atomKK->k_bond_atom,atomKK->bond_atom,nmax,atomKK->bond_per_atom,"atom:bond_atom"); grow_reset(); sync(Host,ALL_MASK); diff --git a/src/KOKKOS/atom_vec_charge_kokkos.cpp b/src/KOKKOS/atom_vec_charge_kokkos.cpp index 7b8b74b4051b63411f673b537ca4c1b9ce731198..a9ae5cc2d1e999d8dc94964e43a3f0d8be4a2a84 100644 --- a/src/KOKKOS/atom_vec_charge_kokkos.cpp +++ b/src/KOKKOS/atom_vec_charge_kokkos.cpp @@ -19,7 +19,7 @@ #include "modify.h" #include "fix.h" #include "atom_masks.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" using namespace LAMMPS_NS; @@ -67,16 +67,16 @@ void AtomVecChargeKokkos::grow(int n) sync(Device,ALL_MASK); modified(Device,ALL_MASK); - memory->grow_kokkos(atomKK->k_tag,atomKK->tag,nmax,"atom:tag"); - memory->grow_kokkos(atomKK->k_type,atomKK->type,nmax,"atom:type"); - memory->grow_kokkos(atomKK->k_mask,atomKK->mask,nmax,"atom:mask"); - memory->grow_kokkos(atomKK->k_image,atomKK->image,nmax,"atom:image"); + memoryKK->grow_kokkos(atomKK->k_tag,atomKK->tag,nmax,"atom:tag"); + memoryKK->grow_kokkos(atomKK->k_type,atomKK->type,nmax,"atom:type"); + memoryKK->grow_kokkos(atomKK->k_mask,atomKK->mask,nmax,"atom:mask"); + memoryKK->grow_kokkos(atomKK->k_image,atomKK->image,nmax,"atom:image"); - memory->grow_kokkos(atomKK->k_x,atomKK->x,nmax,3,"atom:x"); - memory->grow_kokkos(atomKK->k_v,atomKK->v,nmax,3,"atom:v"); - memory->grow_kokkos(atomKK->k_f,atomKK->f,nmax,3,"atom:f"); + memoryKK->grow_kokkos(atomKK->k_x,atomKK->x,nmax,3,"atom:x"); + memoryKK->grow_kokkos(atomKK->k_v,atomKK->v,nmax,3,"atom:v"); + memoryKK->grow_kokkos(atomKK->k_f,atomKK->f,nmax,3,"atom:f"); - memory->grow_kokkos(atomKK->k_q,atomKK->q,nmax,"atom:q"); + memoryKK->grow_kokkos(atomKK->k_q,atomKK->q,nmax,"atom:q"); grow_reset(); sync(Host,ALL_MASK); diff --git a/src/KOKKOS/atom_vec_dpd_kokkos.cpp b/src/KOKKOS/atom_vec_dpd_kokkos.cpp index c4e493bd8536dd036931e61c0e0f080fc9d925ad..9c54ffccc50213111d7e47bd39751f474c475438 100644 --- a/src/KOKKOS/atom_vec_dpd_kokkos.cpp +++ b/src/KOKKOS/atom_vec_dpd_kokkos.cpp @@ -19,7 +19,7 @@ #include "modify.h" #include "fix.h" #include "atom_masks.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" using namespace LAMMPS_NS; @@ -67,24 +67,24 @@ void AtomVecDPDKokkos::grow(int n) sync(Device,ALL_MASK); modified(Device,ALL_MASK); - memory->grow_kokkos(atomKK->k_tag,atomKK->tag,nmax,"atom:tag"); - memory->grow_kokkos(atomKK->k_type,atomKK->type,nmax,"atom:type"); - memory->grow_kokkos(atomKK->k_mask,atomKK->mask,nmax,"atom:mask"); - memory->grow_kokkos(atomKK->k_image,atomKK->image,nmax,"atom:image"); + memoryKK->grow_kokkos(atomKK->k_tag,atomKK->tag,nmax,"atom:tag"); + memoryKK->grow_kokkos(atomKK->k_type,atomKK->type,nmax,"atom:type"); + memoryKK->grow_kokkos(atomKK->k_mask,atomKK->mask,nmax,"atom:mask"); + memoryKK->grow_kokkos(atomKK->k_image,atomKK->image,nmax,"atom:image"); - memory->grow_kokkos(atomKK->k_x,atomKK->x,nmax,3,"atom:x"); - memory->grow_kokkos(atomKK->k_v,atomKK->v,nmax,3,"atom:v"); - memory->grow_kokkos(atomKK->k_f,atomKK->f,nmax,3,"atom:f"); + memoryKK->grow_kokkos(atomKK->k_x,atomKK->x,nmax,3,"atom:x"); + memoryKK->grow_kokkos(atomKK->k_v,atomKK->v,nmax,3,"atom:v"); + memoryKK->grow_kokkos(atomKK->k_f,atomKK->f,nmax,3,"atom:f"); - memory->grow_kokkos(atomKK->k_rho,atomKK->rho,nmax,"atom:rho"); - memory->grow_kokkos(atomKK->k_dpdTheta,atomKK->dpdTheta,nmax,"atom:dpdTheta"); - memory->grow_kokkos(atomKK->k_uCond,atomKK->uCond,nmax,"atom:uCond"); - memory->grow_kokkos(atomKK->k_uMech,atomKK->uMech,nmax,"atom:uMech"); - memory->grow_kokkos(atomKK->k_uChem,atomKK->uChem,nmax,"atom:uChem"); - memory->grow_kokkos(atomKK->k_uCG,atomKK->uCG,nmax,"atom:uCG"); - memory->grow_kokkos(atomKK->k_uCGnew,atomKK->uCGnew,nmax,"atom:uCGnew"); - memory->grow_kokkos(atomKK->k_duChem,atomKK->duChem,nmax,"atom:duChem"); + memoryKK->grow_kokkos(atomKK->k_rho,atomKK->rho,nmax,"atom:rho"); + memoryKK->grow_kokkos(atomKK->k_dpdTheta,atomKK->dpdTheta,nmax,"atom:dpdTheta"); + memoryKK->grow_kokkos(atomKK->k_uCond,atomKK->uCond,nmax,"atom:uCond"); + memoryKK->grow_kokkos(atomKK->k_uMech,atomKK->uMech,nmax,"atom:uMech"); + memoryKK->grow_kokkos(atomKK->k_uChem,atomKK->uChem,nmax,"atom:uChem"); + memoryKK->grow_kokkos(atomKK->k_uCG,atomKK->uCG,nmax,"atom:uCG"); + memoryKK->grow_kokkos(atomKK->k_uCGnew,atomKK->uCGnew,nmax,"atom:uCGnew"); + memoryKK->grow_kokkos(atomKK->k_duChem,atomKK->duChem,nmax,"atom:duChem"); if (atom->nextra_grow) for (int iextra = 0; iextra < atom->nextra_grow; iextra++) diff --git a/src/KOKKOS/atom_vec_full_kokkos.cpp b/src/KOKKOS/atom_vec_full_kokkos.cpp index 8e9abe40675f7816ca969a23c0ab55855f71c5d1..9369d7e84437f184051aed76f129cbb3fe028efb 100644 --- a/src/KOKKOS/atom_vec_full_kokkos.cpp +++ b/src/KOKKOS/atom_vec_full_kokkos.cpp @@ -19,7 +19,7 @@ #include "modify.h" #include "fix.h" #include "atom_masks.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" using namespace LAMMPS_NS; @@ -67,59 +67,59 @@ void AtomVecFullKokkos::grow(int n) sync(Device,ALL_MASK); modified(Device,ALL_MASK); - memory->grow_kokkos(atomKK->k_tag,atomKK->tag,nmax,"atom:tag"); - memory->grow_kokkos(atomKK->k_type,atomKK->type,nmax,"atom:type"); - memory->grow_kokkos(atomKK->k_mask,atomKK->mask,nmax,"atom:mask"); - memory->grow_kokkos(atomKK->k_image,atomKK->image,nmax,"atom:image"); + memoryKK->grow_kokkos(atomKK->k_tag,atomKK->tag,nmax,"atom:tag"); + memoryKK->grow_kokkos(atomKK->k_type,atomKK->type,nmax,"atom:type"); + memoryKK->grow_kokkos(atomKK->k_mask,atomKK->mask,nmax,"atom:mask"); + memoryKK->grow_kokkos(atomKK->k_image,atomKK->image,nmax,"atom:image"); - memory->grow_kokkos(atomKK->k_x,atomKK->x,nmax,3,"atom:x"); - memory->grow_kokkos(atomKK->k_v,atomKK->v,nmax,3,"atom:v"); - memory->grow_kokkos(atomKK->k_f,atomKK->f,nmax,3,"atom:f"); + memoryKK->grow_kokkos(atomKK->k_x,atomKK->x,nmax,3,"atom:x"); + memoryKK->grow_kokkos(atomKK->k_v,atomKK->v,nmax,3,"atom:v"); + memoryKK->grow_kokkos(atomKK->k_f,atomKK->f,nmax,3,"atom:f"); - memory->grow_kokkos(atomKK->k_q,atomKK->q,nmax,"atom:q"); - memory->grow_kokkos(atomKK->k_molecule,atomKK->molecule,nmax,"atom:molecule"); + memoryKK->grow_kokkos(atomKK->k_q,atomKK->q,nmax,"atom:q"); + memoryKK->grow_kokkos(atomKK->k_molecule,atomKK->molecule,nmax,"atom:molecule"); - memory->grow_kokkos(atomKK->k_nspecial,atomKK->nspecial,nmax,3,"atom:nspecial"); - memory->grow_kokkos(atomKK->k_special,atomKK->special,nmax,atomKK->maxspecial, + memoryKK->grow_kokkos(atomKK->k_nspecial,atomKK->nspecial,nmax,3,"atom:nspecial"); + memoryKK->grow_kokkos(atomKK->k_special,atomKK->special,nmax,atomKK->maxspecial, "atom:special"); - memory->grow_kokkos(atomKK->k_num_bond,atomKK->num_bond,nmax,"atom:num_bond"); - memory->grow_kokkos(atomKK->k_bond_type,atomKK->bond_type,nmax,atomKK->bond_per_atom, + memoryKK->grow_kokkos(atomKK->k_num_bond,atomKK->num_bond,nmax,"atom:num_bond"); + memoryKK->grow_kokkos(atomKK->k_bond_type,atomKK->bond_type,nmax,atomKK->bond_per_atom, "atom:bond_type"); - memory->grow_kokkos(atomKK->k_bond_atom,atomKK->bond_atom,nmax,atomKK->bond_per_atom, + memoryKK->grow_kokkos(atomKK->k_bond_atom,atomKK->bond_atom,nmax,atomKK->bond_per_atom, "atom:bond_atom"); - memory->grow_kokkos(atomKK->k_num_angle,atomKK->num_angle,nmax,"atom:num_angle"); - memory->grow_kokkos(atomKK->k_angle_type,atomKK->angle_type,nmax,atomKK->angle_per_atom, + memoryKK->grow_kokkos(atomKK->k_num_angle,atomKK->num_angle,nmax,"atom:num_angle"); + memoryKK->grow_kokkos(atomKK->k_angle_type,atomKK->angle_type,nmax,atomKK->angle_per_atom, "atom:angle_type"); - memory->grow_kokkos(atomKK->k_angle_atom1,atomKK->angle_atom1,nmax,atomKK->angle_per_atom, + memoryKK->grow_kokkos(atomKK->k_angle_atom1,atomKK->angle_atom1,nmax,atomKK->angle_per_atom, "atom:angle_atom1"); - memory->grow_kokkos(atomKK->k_angle_atom2,atomKK->angle_atom2,nmax,atomKK->angle_per_atom, + memoryKK->grow_kokkos(atomKK->k_angle_atom2,atomKK->angle_atom2,nmax,atomKK->angle_per_atom, "atom:angle_atom2"); - memory->grow_kokkos(atomKK->k_angle_atom3,atomKK->angle_atom3,nmax,atomKK->angle_per_atom, + memoryKK->grow_kokkos(atomKK->k_angle_atom3,atomKK->angle_atom3,nmax,atomKK->angle_per_atom, "atom:angle_atom3"); - memory->grow_kokkos(atomKK->k_num_dihedral,atomKK->num_dihedral,nmax,"atom:num_dihedral"); - memory->grow_kokkos(atomKK->k_dihedral_type,atomKK->dihedral_type,nmax, + memoryKK->grow_kokkos(atomKK->k_num_dihedral,atomKK->num_dihedral,nmax,"atom:num_dihedral"); + memoryKK->grow_kokkos(atomKK->k_dihedral_type,atomKK->dihedral_type,nmax, atomKK->dihedral_per_atom,"atom:dihedral_type"); - memory->grow_kokkos(atomKK->k_dihedral_atom1,atomKK->dihedral_atom1,nmax, + memoryKK->grow_kokkos(atomKK->k_dihedral_atom1,atomKK->dihedral_atom1,nmax, atomKK->dihedral_per_atom,"atom:dihedral_atom1"); - memory->grow_kokkos(atomKK->k_dihedral_atom2,atomKK->dihedral_atom2,nmax, + memoryKK->grow_kokkos(atomKK->k_dihedral_atom2,atomKK->dihedral_atom2,nmax, atomKK->dihedral_per_atom,"atom:dihedral_atom2"); - memory->grow_kokkos(atomKK->k_dihedral_atom3,atomKK->dihedral_atom3,nmax, + memoryKK->grow_kokkos(atomKK->k_dihedral_atom3,atomKK->dihedral_atom3,nmax, atomKK->dihedral_per_atom,"atom:dihedral_atom3"); - memory->grow_kokkos(atomKK->k_dihedral_atom4,atomKK->dihedral_atom4,nmax, + memoryKK->grow_kokkos(atomKK->k_dihedral_atom4,atomKK->dihedral_atom4,nmax, atomKK->dihedral_per_atom,"atom:dihedral_atom4"); - memory->grow_kokkos(atomKK->k_num_improper,atomKK->num_improper,nmax,"atom:num_improper"); - memory->grow_kokkos(atomKK->k_improper_type,atomKK->improper_type,nmax, + memoryKK->grow_kokkos(atomKK->k_num_improper,atomKK->num_improper,nmax,"atom:num_improper"); + memoryKK->grow_kokkos(atomKK->k_improper_type,atomKK->improper_type,nmax, atomKK->improper_per_atom,"atom:improper_type"); - memory->grow_kokkos(atomKK->k_improper_atom1,atomKK->improper_atom1,nmax, + memoryKK->grow_kokkos(atomKK->k_improper_atom1,atomKK->improper_atom1,nmax, atomKK->improper_per_atom,"atom:improper_atom1"); - memory->grow_kokkos(atomKK->k_improper_atom2,atomKK->improper_atom2,nmax, + memoryKK->grow_kokkos(atomKK->k_improper_atom2,atomKK->improper_atom2,nmax, atomKK->improper_per_atom,"atom:improper_atom2"); - memory->grow_kokkos(atomKK->k_improper_atom3,atomKK->improper_atom3,nmax, + memoryKK->grow_kokkos(atomKK->k_improper_atom3,atomKK->improper_atom3,nmax, atomKK->improper_per_atom,"atom:improper_atom3"); - memory->grow_kokkos(atomKK->k_improper_atom4,atomKK->improper_atom4,nmax, + memoryKK->grow_kokkos(atomKK->k_improper_atom4,atomKK->improper_atom4,nmax, atomKK->improper_per_atom,"atom:improper_atom4"); grow_reset(); diff --git a/src/KOKKOS/atom_vec_hybrid_kokkos.cpp b/src/KOKKOS/atom_vec_hybrid_kokkos.cpp index e5e361e70ab21645b67608f6d4f68ab07fb19900..b5aadb18d612cfbd30a6c3666fd85a3dfe489d82 100644 --- a/src/KOKKOS/atom_vec_hybrid_kokkos.cpp +++ b/src/KOKKOS/atom_vec_hybrid_kokkos.cpp @@ -18,7 +18,7 @@ #include "domain.h" #include "modify.h" #include "fix.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" diff --git a/src/KOKKOS/atom_vec_molecular_kokkos.cpp b/src/KOKKOS/atom_vec_molecular_kokkos.cpp index dbf6a857b2521874ba7dca32a4cf374030c546e1..6f232a319b06238c95ff903706ad759228045c82 100644 --- a/src/KOKKOS/atom_vec_molecular_kokkos.cpp +++ b/src/KOKKOS/atom_vec_molecular_kokkos.cpp @@ -19,7 +19,7 @@ #include "modify.h" #include "fix.h" #include "atom_masks.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" using namespace LAMMPS_NS; @@ -67,57 +67,57 @@ void AtomVecMolecularKokkos::grow(int n) sync(Device,ALL_MASK); modified(Device,ALL_MASK); - memory->grow_kokkos(atomKK->k_tag,atomKK->tag,nmax,"atom:tag"); - memory->grow_kokkos(atomKK->k_type,atomKK->type,nmax,"atom:type"); - memory->grow_kokkos(atomKK->k_mask,atomKK->mask,nmax,"atom:mask"); - memory->grow_kokkos(atomKK->k_image,atomKK->image,nmax,"atom:image"); + memoryKK->grow_kokkos(atomKK->k_tag,atomKK->tag,nmax,"atom:tag"); + memoryKK->grow_kokkos(atomKK->k_type,atomKK->type,nmax,"atom:type"); + memoryKK->grow_kokkos(atomKK->k_mask,atomKK->mask,nmax,"atom:mask"); + memoryKK->grow_kokkos(atomKK->k_image,atomKK->image,nmax,"atom:image"); - memory->grow_kokkos(atomKK->k_x,atomKK->x,nmax,3,"atom:x"); - memory->grow_kokkos(atomKK->k_v,atomKK->v,nmax,3,"atom:v"); - memory->grow_kokkos(atomKK->k_f,atomKK->f,nmax,3,"atom:f"); + memoryKK->grow_kokkos(atomKK->k_x,atomKK->x,nmax,3,"atom:x"); + memoryKK->grow_kokkos(atomKK->k_v,atomKK->v,nmax,3,"atom:v"); + memoryKK->grow_kokkos(atomKK->k_f,atomKK->f,nmax,3,"atom:f"); - memory->grow_kokkos(atomKK->k_molecule,atomKK->molecule,nmax,"atom:molecule"); - memory->grow_kokkos(atomKK->k_nspecial,atomKK->nspecial,nmax,3,"atom:nspecial"); - memory->grow_kokkos(atomKK->k_special,atomKK->special,nmax,atomKK->maxspecial, + memoryKK->grow_kokkos(atomKK->k_molecule,atomKK->molecule,nmax,"atom:molecule"); + memoryKK->grow_kokkos(atomKK->k_nspecial,atomKK->nspecial,nmax,3,"atom:nspecial"); + memoryKK->grow_kokkos(atomKK->k_special,atomKK->special,nmax,atomKK->maxspecial, "atom:special"); - memory->grow_kokkos(atomKK->k_num_bond,atomKK->num_bond,nmax,"atom:num_bond"); - memory->grow_kokkos(atomKK->k_bond_type,atomKK->bond_type,nmax,atomKK->bond_per_atom, + memoryKK->grow_kokkos(atomKK->k_num_bond,atomKK->num_bond,nmax,"atom:num_bond"); + memoryKK->grow_kokkos(atomKK->k_bond_type,atomKK->bond_type,nmax,atomKK->bond_per_atom, "atom:bond_type"); - memory->grow_kokkos(atomKK->k_bond_atom,atomKK->bond_atom,nmax,atomKK->bond_per_atom, + memoryKK->grow_kokkos(atomKK->k_bond_atom,atomKK->bond_atom,nmax,atomKK->bond_per_atom, "atom:bond_atom"); - memory->grow_kokkos(atomKK->k_num_angle,atomKK->num_angle,nmax,"atom:num_angle"); - memory->grow_kokkos(atomKK->k_angle_type,atomKK->angle_type,nmax,atomKK->angle_per_atom, + memoryKK->grow_kokkos(atomKK->k_num_angle,atomKK->num_angle,nmax,"atom:num_angle"); + memoryKK->grow_kokkos(atomKK->k_angle_type,atomKK->angle_type,nmax,atomKK->angle_per_atom, "atom:angle_type"); - memory->grow_kokkos(atomKK->k_angle_atom1,atomKK->angle_atom1,nmax,atomKK->angle_per_atom, + memoryKK->grow_kokkos(atomKK->k_angle_atom1,atomKK->angle_atom1,nmax,atomKK->angle_per_atom, "atom:angle_atom1"); - memory->grow_kokkos(atomKK->k_angle_atom2,atomKK->angle_atom2,nmax,atomKK->angle_per_atom, + memoryKK->grow_kokkos(atomKK->k_angle_atom2,atomKK->angle_atom2,nmax,atomKK->angle_per_atom, "atom:angle_atom2"); - memory->grow_kokkos(atomKK->k_angle_atom3,atomKK->angle_atom3,nmax,atomKK->angle_per_atom, + memoryKK->grow_kokkos(atomKK->k_angle_atom3,atomKK->angle_atom3,nmax,atomKK->angle_per_atom, "atom:angle_atom3"); - memory->grow_kokkos(atomKK->k_num_dihedral,atomKK->num_dihedral,nmax,"atom:num_dihedral"); - memory->grow_kokkos(atomKK->k_dihedral_type,atomKK->dihedral_type,nmax, + memoryKK->grow_kokkos(atomKK->k_num_dihedral,atomKK->num_dihedral,nmax,"atom:num_dihedral"); + memoryKK->grow_kokkos(atomKK->k_dihedral_type,atomKK->dihedral_type,nmax, atomKK->dihedral_per_atom,"atom:dihedral_type"); - memory->grow_kokkos(atomKK->k_dihedral_atom1,atomKK->dihedral_atom1,nmax, + memoryKK->grow_kokkos(atomKK->k_dihedral_atom1,atomKK->dihedral_atom1,nmax, atomKK->dihedral_per_atom,"atom:dihedral_atom1"); - memory->grow_kokkos(atomKK->k_dihedral_atom2,atomKK->dihedral_atom2,nmax, + memoryKK->grow_kokkos(atomKK->k_dihedral_atom2,atomKK->dihedral_atom2,nmax, atomKK->dihedral_per_atom,"atom:dihedral_atom2"); - memory->grow_kokkos(atomKK->k_dihedral_atom3,atomKK->dihedral_atom3,nmax, + memoryKK->grow_kokkos(atomKK->k_dihedral_atom3,atomKK->dihedral_atom3,nmax, atomKK->dihedral_per_atom,"atom:dihedral_atom3"); - memory->grow_kokkos(atomKK->k_dihedral_atom4,atomKK->dihedral_atom4,nmax, + memoryKK->grow_kokkos(atomKK->k_dihedral_atom4,atomKK->dihedral_atom4,nmax, atomKK->dihedral_per_atom,"atom:dihedral_atom4"); - memory->grow_kokkos(atomKK->k_num_improper,atomKK->num_improper,nmax,"atom:num_improper"); - memory->grow_kokkos(atomKK->k_improper_type,atomKK->improper_type,nmax, + memoryKK->grow_kokkos(atomKK->k_num_improper,atomKK->num_improper,nmax,"atom:num_improper"); + memoryKK->grow_kokkos(atomKK->k_improper_type,atomKK->improper_type,nmax, atomKK->improper_per_atom,"atom:improper_type"); - memory->grow_kokkos(atomKK->k_improper_atom1,atomKK->improper_atom1,nmax, + memoryKK->grow_kokkos(atomKK->k_improper_atom1,atomKK->improper_atom1,nmax, atomKK->improper_per_atom,"atom:improper_atom1"); - memory->grow_kokkos(atomKK->k_improper_atom2,atomKK->improper_atom2,nmax, + memoryKK->grow_kokkos(atomKK->k_improper_atom2,atomKK->improper_atom2,nmax, atomKK->improper_per_atom,"atom:improper_atom2"); - memory->grow_kokkos(atomKK->k_improper_atom3,atomKK->improper_atom3,nmax, + memoryKK->grow_kokkos(atomKK->k_improper_atom3,atomKK->improper_atom3,nmax, atomKK->improper_per_atom,"atom:improper_atom3"); - memory->grow_kokkos(atomKK->k_improper_atom4,atomKK->improper_atom4,nmax, + memoryKK->grow_kokkos(atomKK->k_improper_atom4,atomKK->improper_atom4,nmax, atomKK->improper_per_atom,"atom:improper_atom4"); grow_reset(); diff --git a/src/KOKKOS/bond_class2_kokkos.cpp b/src/KOKKOS/bond_class2_kokkos.cpp index b3c11c9a06c468be08f9ce2b2a29919c949ff475..df2f2c1e9be24a6d807ecf12e67c239735dde7f6 100644 --- a/src/KOKKOS/bond_class2_kokkos.cpp +++ b/src/KOKKOS/bond_class2_kokkos.cpp @@ -23,7 +23,7 @@ #include "domain.h" #include "comm.h" #include "force.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -47,8 +47,8 @@ template<class DeviceType> BondClass2Kokkos<DeviceType>::~BondClass2Kokkos() { if (!copymode) { - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); } } @@ -67,15 +67,15 @@ void BondClass2Kokkos<DeviceType>::compute(int eflag_in, int vflag_in) if (eflag_atom) { //if(k_eatom.dimension_0()<maxeatom) { // won't work without adding zero functor - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom"); d_eatom = k_eatom.template view<DeviceType>(); //} } if (vflag_atom) { //if(k_vatom.dimension_0()<maxvatom) { // won't work without adding zero functor - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"improper:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"improper:vatom"); d_vatom = k_vatom.template view<DeviceType>(); //} } diff --git a/src/KOKKOS/bond_fene_kokkos.cpp b/src/KOKKOS/bond_fene_kokkos.cpp index 8a716a98ef534bfdcc3daf31696d01e1af22b071..20c20542083b35c289d8efdab2f3dce5adc95077 100644 --- a/src/KOKKOS/bond_fene_kokkos.cpp +++ b/src/KOKKOS/bond_fene_kokkos.cpp @@ -23,7 +23,7 @@ #include "domain.h" #include "comm.h" #include "force.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -56,8 +56,8 @@ template<class DeviceType> BondFENEKokkos<DeviceType>::~BondFENEKokkos() { if (!copymode) { - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); } } @@ -75,13 +75,13 @@ void BondFENEKokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"bond:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"bond:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"bond:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"bond:vatom"); d_vatom = k_vatom.view<DeviceType>(); } diff --git a/src/KOKKOS/bond_harmonic_kokkos.cpp b/src/KOKKOS/bond_harmonic_kokkos.cpp index da45c70d6c7739f08ad0fb559e6d9d7f768437c4..c4e0c3a81763e765ee593ba13b84c4b0fcec3616 100644 --- a/src/KOKKOS/bond_harmonic_kokkos.cpp +++ b/src/KOKKOS/bond_harmonic_kokkos.cpp @@ -23,7 +23,7 @@ #include "domain.h" #include "comm.h" #include "force.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -48,8 +48,8 @@ template<class DeviceType> BondHarmonicKokkos<DeviceType>::~BondHarmonicKokkos() { if (!copymode) { - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); } } @@ -68,15 +68,15 @@ void BondHarmonicKokkos<DeviceType>::compute(int eflag_in, int vflag_in) if (eflag_atom) { //if(k_eatom.dimension_0()<maxeatom) { // won't work without adding zero functor - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom"); d_eatom = k_eatom.template view<DeviceType>(); //} } if (vflag_atom) { //if(k_vatom.dimension_0()<maxvatom) { // won't work without adding zero functor - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"improper:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"improper:vatom"); d_vatom = k_vatom.template view<DeviceType>(); //} } diff --git a/src/KOKKOS/comm_kokkos.cpp b/src/KOKKOS/comm_kokkos.cpp index 5534341342d5017378169609dce1a0c5964d8338..e506fa1ad42c849d0b31b0e564a45842d6aa1cfe 100644 --- a/src/KOKKOS/comm_kokkos.cpp +++ b/src/KOKKOS/comm_kokkos.cpp @@ -20,7 +20,7 @@ #include "domain.h" #include "atom_masks.h" #include "error.h" -#include "memory.h" +#include "memory_kokkos.h" #include "force.h" #include "pair.h" #include "fix.h" @@ -28,6 +28,7 @@ #include "dump.h" #include "output.h" #include "modify.h" +#include "kokkos_base.h" using namespace LAMMPS_NS; @@ -71,7 +72,7 @@ CommKokkos::CommKokkos(LAMMPS *lmp) : CommBrick(lmp) for (int i = 0; i < maxswap; i++) { maxsendlist[i] = BUFMIN; } - memory->create_kokkos(k_sendlist,sendlist,maxswap,BUFMIN,"comm:sendlist"); + memoryKK->create_kokkos(k_sendlist,sendlist,maxswap,BUFMIN,"comm:sendlist"); max_buf_pair = 0; k_buf_send_pair = DAT::tdual_xfloat_1d("comm:k_buf_send_pair",1); @@ -82,11 +83,11 @@ CommKokkos::CommKokkos(LAMMPS *lmp) : CommBrick(lmp) CommKokkos::~CommKokkos() { - memory->destroy_kokkos(k_sendlist,sendlist); + memoryKK->destroy_kokkos(k_sendlist,sendlist); sendlist = NULL; - memory->destroy_kokkos(k_buf_send,buf_send); + memoryKK->destroy_kokkos(k_buf_send,buf_send); buf_send = NULL; - memory->destroy_kokkos(k_buf_recv,buf_recv); + memoryKK->destroy_kokkos(k_buf_recv,buf_recv); buf_recv = NULL; } @@ -379,6 +380,7 @@ void CommKokkos::forward_comm_pair_device(Pair *pair) MPI_Request request; int nsize = pair->comm_forward; + KokkosBase* pairKKBase = dynamic_cast<KokkosBase*>(pair); for (iswap = 0; iswap < nswap; iswap++) { int n = MAX(max_buf_pair,nsize*sendnum[iswap]); @@ -391,7 +393,7 @@ void CommKokkos::forward_comm_pair_device(Pair *pair) // pack buffer - n = pair->pack_forward_comm_kokkos(sendnum[iswap],k_sendlist, + n = pairKKBase->pack_forward_comm_kokkos(sendnum[iswap],k_sendlist, iswap,k_buf_send_pair,pbc_flag[iswap],pbc[iswap]); // exchange with another proc @@ -408,7 +410,7 @@ void CommKokkos::forward_comm_pair_device(Pair *pair) // unpack buffer - pair->unpack_forward_comm_kokkos(recvnum[iswap],firstrecv[iswap],k_buf_recv_pair); + pairKKBase->unpack_forward_comm_kokkos(recvnum[iswap],firstrecv[iswap],k_buf_recv_pair); } } @@ -1067,7 +1069,7 @@ void CommKokkos::grow_list(int iswap, int n) k_sendlist.modify<LMPHostType>(); } - memory->grow_kokkos(k_sendlist,sendlist,maxswap,size,"comm:sendlist"); + memoryKK->grow_kokkos(k_sendlist,sendlist,maxswap,size,"comm:sendlist"); for(int i=0;i<maxswap;i++) { maxsendlist[i]=size; sendlist[i]=&k_sendlist.view<LMPHostType>()(i,0); @@ -1095,7 +1097,7 @@ void CommKokkos::grow_swap(int n) k_sendlist.modify<LMPHostType>(); } - memory->grow_kokkos(k_sendlist,sendlist,maxswap,size,"comm:sendlist"); + memoryKK->grow_kokkos(k_sendlist,sendlist,maxswap,size,"comm:sendlist"); memory->grow(maxsendlist,n,"comm:maxsendlist"); for (int i=0;i<maxswap;i++) maxsendlist[i]=size; diff --git a/src/KOKKOS/comm_tiled_kokkos.cpp b/src/KOKKOS/comm_tiled_kokkos.cpp index adcc634aa1d766858c62e85a9ec6f7557d37da6c..33cd8eaa6e9eabc518a163d1ffbe68bbe0cffc30 100644 --- a/src/KOKKOS/comm_tiled_kokkos.cpp +++ b/src/KOKKOS/comm_tiled_kokkos.cpp @@ -25,7 +25,7 @@ #include "compute.h" #include "output.h" #include "dump.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" diff --git a/src/KOKKOS/dihedral_charmm_kokkos.cpp b/src/KOKKOS/dihedral_charmm_kokkos.cpp index 7f2117c97f545021ca35172c62f9299adb9e8275..71635ec76c48d5642e7eccccef8d360c2077fc60 100644 --- a/src/KOKKOS/dihedral_charmm_kokkos.cpp +++ b/src/KOKKOS/dihedral_charmm_kokkos.cpp @@ -25,7 +25,7 @@ #include "force.h" #include "pair.h" #include "math_const.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -56,8 +56,8 @@ template<class DeviceType> DihedralCharmmKokkos<DeviceType>::~DihedralCharmmKokkos() { if (!copymode) { - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); } } @@ -81,8 +81,8 @@ void DihedralCharmmKokkos<DeviceType>::compute(int eflag_in, int vflag_in) if (eflag_atom) { //if(k_eatom.dimension_0()<maxeatom) { // won't work without adding zero functor - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"dihedral:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"dihedral:eatom"); d_eatom = k_eatom.template view<DeviceType>(); k_eatom_pair = Kokkos::DualView<E_FLOAT*,Kokkos::LayoutRight,DeviceType>("dihedral:eatom_pair",maxeatom); d_eatom_pair = k_eatom.template view<DeviceType>(); @@ -90,8 +90,8 @@ void DihedralCharmmKokkos<DeviceType>::compute(int eflag_in, int vflag_in) } if (vflag_atom) { //if(k_vatom.dimension_0()<maxvatom) { // won't work without adding zero functor - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"dihedral:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"dihedral:vatom"); d_vatom = k_vatom.template view<DeviceType>(); k_vatom_pair = Kokkos::DualView<F_FLOAT*[6],Kokkos::LayoutRight,DeviceType>("dihedral:vatom_pair",maxvatom); d_vatom_pair = k_vatom.template view<DeviceType>(); diff --git a/src/KOKKOS/dihedral_class2_kokkos.cpp b/src/KOKKOS/dihedral_class2_kokkos.cpp index 89e42c6f836a1f3fa9e502875332bffb8bb4cd08..d32ea4a4611089808f31faf7756967f6572e0f94 100644 --- a/src/KOKKOS/dihedral_class2_kokkos.cpp +++ b/src/KOKKOS/dihedral_class2_kokkos.cpp @@ -24,7 +24,7 @@ #include "domain.h" #include "force.h" #include "update.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -56,8 +56,8 @@ template<class DeviceType> DihedralClass2Kokkos<DeviceType>::~DihedralClass2Kokkos() { if (!copymode) { - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); } } @@ -75,13 +75,13 @@ void DihedralClass2Kokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"dihedral:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"dihedral:eatom"); d_eatom = k_eatom.template view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"dihedral:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"dihedral:vatom"); d_vatom = k_vatom.template view<DeviceType>(); } diff --git a/src/KOKKOS/dihedral_opls_kokkos.cpp b/src/KOKKOS/dihedral_opls_kokkos.cpp index 0ee00ca8db438c2c47eba7443e93039295ab2f93..4349aff48df93848f097bd1130c758ec6be91951 100644 --- a/src/KOKKOS/dihedral_opls_kokkos.cpp +++ b/src/KOKKOS/dihedral_opls_kokkos.cpp @@ -24,7 +24,7 @@ #include "domain.h" #include "force.h" #include "update.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -56,8 +56,8 @@ template<class DeviceType> DihedralOPLSKokkos<DeviceType>::~DihedralOPLSKokkos() { if (!copymode) { - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); } } @@ -75,13 +75,13 @@ void DihedralOPLSKokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"dihedral:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"dihedral:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"dihedral:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"dihedral:vatom"); d_vatom = k_vatom.view<DeviceType>(); } diff --git a/src/KOKKOS/fix_eos_table_rx_kokkos.cpp b/src/KOKKOS/fix_eos_table_rx_kokkos.cpp index 552141ced24f50d68108d372405c5cb82da03360..5c106c19f3432edcc5a432d96031b12270bb39bb 100644 --- a/src/KOKKOS/fix_eos_table_rx_kokkos.cpp +++ b/src/KOKKOS/fix_eos_table_rx_kokkos.cpp @@ -21,7 +21,7 @@ #include "atom_kokkos.h" #include "error.h" #include "force.h" -#include "memory.h" +#include "memory_kokkos.h" #include "comm.h" #include <math.h> #include "modify.h" @@ -517,14 +517,14 @@ void FixEOStableRXKokkos<DeviceType>::create_kokkos_tables() { const int tlm1 = tablength-1; - memory->create_kokkos(d_table->lo,h_table->lo,ntables,"Table::lo"); - memory->create_kokkos(d_table->hi,h_table->hi,ntables,"Table::hi"); - memory->create_kokkos(d_table->invdelta,h_table->invdelta,ntables,"Table::invdelta"); + memoryKK->create_kokkos(d_table->lo,h_table->lo,ntables,"Table::lo"); + memoryKK->create_kokkos(d_table->hi,h_table->hi,ntables,"Table::hi"); + memoryKK->create_kokkos(d_table->invdelta,h_table->invdelta,ntables,"Table::invdelta"); if(tabstyle == LINEAR) { - memory->create_kokkos(d_table->r,h_table->r,ntables,tablength,"Table::r"); - memory->create_kokkos(d_table->e,h_table->e,ntables,tablength,"Table::e"); - memory->create_kokkos(d_table->de,h_table->de,ntables,tlm1,"Table::de"); + memoryKK->create_kokkos(d_table->r,h_table->r,ntables,tablength,"Table::r"); + memoryKK->create_kokkos(d_table->e,h_table->e,ntables,tablength,"Table::e"); + memoryKK->create_kokkos(d_table->de,h_table->de,ntables,tlm1,"Table::de"); } for(int i=0; i < ntables; i++) { diff --git a/src/KOKKOS/fix_langevin_kokkos.cpp b/src/KOKKOS/fix_langevin_kokkos.cpp index fb0f329a91bb398bad243493e9d834ac3ab1aa8d..108c3b692ad7f4fcb176a407a558c5dca3685ca8 100644 --- a/src/KOKKOS/fix_langevin_kokkos.cpp +++ b/src/KOKKOS/fix_langevin_kokkos.cpp @@ -21,7 +21,7 @@ #include "update.h" #include "respa.h" #include "error.h" -#include "memory.h" +#include "memory_kokkos.h" #include "group.h" #include "random_mars.h" #include "compute.h" @@ -49,9 +49,9 @@ FixLangevinKokkos<DeviceType>::FixLangevinKokkos(LAMMPS *lmp, int narg, char **a int ntypes = atomKK->ntypes; // allocate per-type arrays for force prefactors - memory->create_kokkos(k_gfactor1,gfactor1,ntypes+1,"langevin:gfactor1"); - memory->create_kokkos(k_gfactor2,gfactor2,ntypes+1,"langevin:gfactor2"); - memory->create_kokkos(k_ratio,ratio,ntypes+1,"langevin:ratio"); + memoryKK->create_kokkos(k_gfactor1,gfactor1,ntypes+1,"langevin:gfactor1"); + memoryKK->create_kokkos(k_gfactor2,gfactor2,ntypes+1,"langevin:gfactor2"); + memoryKK->create_kokkos(k_ratio,ratio,ntypes+1,"langevin:ratio"); d_gfactor1 = k_gfactor1.template view<DeviceType>(); h_gfactor1 = k_gfactor1.template view<LMPHostType>(); d_gfactor2 = k_gfactor2.template view<DeviceType>(); @@ -92,12 +92,12 @@ FixLangevinKokkos<DeviceType>::FixLangevinKokkos(LAMMPS *lmp, int narg, char **a template<class DeviceType> FixLangevinKokkos<DeviceType>::~FixLangevinKokkos() { - memory->destroy_kokkos(k_gfactor1,gfactor1); - memory->destroy_kokkos(k_gfactor2,gfactor2); - memory->destroy_kokkos(k_ratio,ratio); - memory->destroy_kokkos(k_flangevin,flangevin); - if(gjfflag) memory->destroy_kokkos(k_franprev,franprev); - memory->destroy_kokkos(k_tforce,tforce); + memoryKK->destroy_kokkos(k_gfactor1,gfactor1); + memoryKK->destroy_kokkos(k_gfactor2,gfactor2); + memoryKK->destroy_kokkos(k_ratio,ratio); + memoryKK->destroy_kokkos(k_flangevin,flangevin); + if(gjfflag) memoryKK->destroy_kokkos(k_franprev,franprev); + memoryKK->destroy_kokkos(k_tforce,tforce); } /* ---------------------------------------------------------------------- */ @@ -121,7 +121,7 @@ void FixLangevinKokkos<DeviceType>::init() template<class DeviceType> void FixLangevinKokkos<DeviceType>::grow_arrays(int nmax) { - memory->grow_kokkos(k_franprev,franprev,nmax,3,"langevin:franprev"); + memoryKK->grow_kokkos(k_franprev,franprev,nmax,3,"langevin:franprev"); d_franprev = k_franprev.template view<DeviceType>(); h_franprev = k_franprev.template view<LMPHostType>(); } @@ -167,9 +167,9 @@ void FixLangevinKokkos<DeviceType>::post_force(int vflag) // reallocate flangevin if necessary if (tallyflag) { if (nlocal > maxatom1) { - memory->destroy_kokkos(k_flangevin,flangevin); + memoryKK->destroy_kokkos(k_flangevin,flangevin); maxatom1 = atomKK->nmax; - memory->create_kokkos(k_flangevin,flangevin,maxatom1,3,"langevin:flangevin"); + memoryKK->create_kokkos(k_flangevin,flangevin,maxatom1,3,"langevin:flangevin"); d_flangevin = k_flangevin.template view<DeviceType>(); h_flangevin = k_flangevin.template view<LMPHostType>(); } @@ -671,8 +671,8 @@ void FixLangevinKokkos<DeviceType>::compute_target() } else { if (atom->nmax > maxatom2) { maxatom2 = atom->nmax; - memory->destroy_kokkos(k_tforce,tforce); - memory->create_kokkos(k_tforce,tforce,maxatom2,"langevin:tforce"); + memoryKK->destroy_kokkos(k_tforce,tforce); + memoryKK->create_kokkos(k_tforce,tforce,maxatom2,"langevin:tforce"); d_tforce = k_tforce.template view<DeviceType>(); h_tforce = k_tforce.template view<LMPHostType>(); } diff --git a/src/KOKKOS/fix_nh_kokkos.cpp b/src/KOKKOS/fix_nh_kokkos.cpp index 7136c776a1de7e39d670fa53c69945f8f1d15b3f..345259e35519b24a6b6b669c175fc9e5bdeaf746 100644 --- a/src/KOKKOS/fix_nh_kokkos.cpp +++ b/src/KOKKOS/fix_nh_kokkos.cpp @@ -33,7 +33,7 @@ #include "update.h" #include "respa.h" #include "domain_kokkos.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" #include "atom_kokkos.h" diff --git a/src/KOKKOS/fix_property_atom_kokkos.cpp b/src/KOKKOS/fix_property_atom_kokkos.cpp index cb52988c318fcbd147a534ac16d2ea5494f5f63d..fe2f101e561f1dca37852084d6be56a6896f8b2d 100644 --- a/src/KOKKOS/fix_property_atom_kokkos.cpp +++ b/src/KOKKOS/fix_property_atom_kokkos.cpp @@ -16,7 +16,7 @@ #include "fix_property_atom_kokkos.h" #include "atom_kokkos.h" #include "comm.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "update.h" @@ -60,7 +60,7 @@ void FixPropertyAtomKokkos::grow_arrays(int nmax) size_t nbytes = (nmax-nmax_old) * sizeof(int); memset(&atom->ivector[index[m]][nmax_old],0,nbytes); } else if (style[m] == DOUBLE) { - memory->grow_kokkos(atomKK->k_dvector,atomKK->dvector,atomKK->k_dvector.dimension_0(),nmax, + memoryKK->grow_kokkos(atomKK->k_dvector,atomKK->dvector,atomKK->k_dvector.dimension_0(),nmax, "atom:dvector"); //memory->grow(atom->dvector[index[m]],nmax,"atom:dvector"); //size_t nbytes = (nmax-nmax_old) * sizeof(double); diff --git a/src/KOKKOS/fix_qeq_reax_kokkos.cpp b/src/KOKKOS/fix_qeq_reax_kokkos.cpp index 5d2f6a0438a400302785c70975beee07d4d4f32b..91a22361fc91b4371eb7c7638e4fdb1b9c229d46 100644 --- a/src/KOKKOS/fix_qeq_reax_kokkos.cpp +++ b/src/KOKKOS/fix_qeq_reax_kokkos.cpp @@ -35,7 +35,7 @@ #include "integrate.h" #include "respa.h" #include "math_const.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "pair_reaxc_kokkos.h" diff --git a/src/KOKKOS/fix_reaxc_bonds_kokkos.cpp b/src/KOKKOS/fix_reaxc_bonds_kokkos.cpp index 0d74a49ed304498a5f11d55133898a4565e56002..586daadd55e64c3922f52ee16faa454f7b723da5 100644 --- a/src/KOKKOS/fix_reaxc_bonds_kokkos.cpp +++ b/src/KOKKOS/fix_reaxc_bonds_kokkos.cpp @@ -31,7 +31,7 @@ #include "compute.h" #include "input.h" #include "variable.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "reaxc_list.h" #include "reaxc_types.h" @@ -95,7 +95,7 @@ void FixReaxCBondsKokkos::Output_ReaxC_Bonds(bigint ntimestep, FILE *fp) MPI_Allreduce(&nlocal,&nlocal_max,1,MPI_INT,MPI_MAX,world); nbuf = 1+(numbonds_max*2+10)*nlocal_max; - memory->create_kokkos(k_buf,buf,nbuf,"reax/c/bonds:buf"); + memoryKK->create_kokkos(k_buf,buf,nbuf,"reax/c/bonds:buf"); // Pass information to buffer if (reaxc->execution_space == Device) @@ -107,7 +107,7 @@ void FixReaxCBondsKokkos::Output_ReaxC_Bonds(bigint ntimestep, FILE *fp) // Receive information from buffer for output RecvBuffer(buf, nbuf, nbuf_local, nlocal_tot, numbonds_max); - memory->destroy_kokkos(k_buf,buf); + memoryKK->destroy_kokkos(k_buf,buf); } /* ---------------------------------------------------------------------- */ diff --git a/src/KOKKOS/fix_reaxc_species_kokkos.cpp b/src/KOKKOS/fix_reaxc_species_kokkos.cpp index f2719f9f0e7048a6db47bf49dfd09979a6f5074b..a676c7ef27ae81cf7ec39e8467e7ecdceb816e06 100644 --- a/src/KOKKOS/fix_reaxc_species_kokkos.cpp +++ b/src/KOKKOS/fix_reaxc_species_kokkos.cpp @@ -33,7 +33,7 @@ #include "compute.h" #include "input.h" #include "variable.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "reaxc_list.h" #include "atom_masks.h" diff --git a/src/KOKKOS/fix_rx_kokkos.cpp b/src/KOKKOS/fix_rx_kokkos.cpp index b1cfd20be23f57a25981e151a5098c0b43b13ae6..61290ece33dfb72c50dc787542e4cea50fe5a077 100644 --- a/src/KOKKOS/fix_rx_kokkos.cpp +++ b/src/KOKKOS/fix_rx_kokkos.cpp @@ -17,7 +17,7 @@ #include "atom_masks.h" #include "atom_kokkos.h" #include "force.h" -#include "memory.h" +#include "memory_kokkos.h" #include "update.h" #include "respa.h" #include "modify.h" @@ -26,6 +26,9 @@ #include "neigh_request.h" #include "error.h" #include "math_special_kokkos.h" +#include "comm.h" +#include "domain.h" +#include "kokkos.h" #include <float.h> // DBL_EPSILON @@ -81,15 +84,15 @@ FixRxKokkos<DeviceType>::~FixRxKokkos() if (copymode) return; if (localTempFlag) - memory->destroy_kokkos(k_dpdThetaLocal, dpdThetaLocal); + memoryKK->destroy_kokkos(k_dpdThetaLocal, dpdThetaLocal); - memory->destroy_kokkos(k_sumWeights, sumWeights); - //memory->destroy_kokkos(k_sumWeights); + memoryKK->destroy_kokkos(k_sumWeights, sumWeights); + //memoryKK->destroy_kokkos(k_sumWeights); //delete [] scratchSpace; - memory->destroy_kokkos(d_scratchSpace); + memoryKK->destroy_kokkos(d_scratchSpace); - memory->destroy_kokkos(k_cutsq); + memoryKK->destroy_kokkos(k_cutsq); } /* ---------------------------------------------------------------------- */ @@ -1233,9 +1236,9 @@ void FixRxKokkos<DeviceType>::create_kinetics_data(void) { //printf("Inside FixRxKokkos::create_kinetics_data\n"); - memory->create_kokkos( d_kineticsData.Arr, h_kineticsData.Arr, nreactions, "KineticsType::Arr"); - memory->create_kokkos( d_kineticsData.nArr, h_kineticsData.nArr, nreactions, "KineticsType::nArr"); - memory->create_kokkos( d_kineticsData.Ea, h_kineticsData.Ea, nreactions, "KineticsType::Ea"); + memoryKK->create_kokkos( d_kineticsData.Arr, h_kineticsData.Arr, nreactions, "KineticsType::Arr"); + memoryKK->create_kokkos( d_kineticsData.nArr, h_kineticsData.nArr, nreactions, "KineticsType::nArr"); + memoryKK->create_kokkos( d_kineticsData.Ea, h_kineticsData.Ea, nreactions, "KineticsType::Ea"); for (int i = 0; i < nreactions; ++i) { @@ -1251,8 +1254,8 @@ void FixRxKokkos<DeviceType>::create_kinetics_data(void) if (useSparseKinetics) { - memory->create_kokkos( d_kineticsData.nu , h_kineticsData.nu , nreactions, sparseKinetics_maxSpecies, "KineticsType::nu"); - memory->create_kokkos( d_kineticsData.nuk, h_kineticsData.nuk, nreactions, sparseKinetics_maxSpecies, "KineticsType::nuk"); + memoryKK->create_kokkos( d_kineticsData.nu , h_kineticsData.nu , nreactions, sparseKinetics_maxSpecies, "KineticsType::nu"); + memoryKK->create_kokkos( d_kineticsData.nuk, h_kineticsData.nuk, nreactions, sparseKinetics_maxSpecies, "KineticsType::nuk"); for (int i = 0; i < nreactions; ++i) for (int k = 0; k < sparseKinetics_maxSpecies; ++k) @@ -1266,8 +1269,8 @@ void FixRxKokkos<DeviceType>::create_kinetics_data(void) if (SparseKinetics_enableIntegralReactions) { - memory->create_kokkos( d_kineticsData.inu, h_kineticsData.inu, nreactions, sparseKinetics_maxSpecies, "KineticsType::inu"); - memory->create_kokkos( d_kineticsData.isIntegral, h_kineticsData.isIntegral, nreactions, "KineticsType::isIntegral"); + memoryKK->create_kokkos( d_kineticsData.inu, h_kineticsData.inu, nreactions, sparseKinetics_maxSpecies, "KineticsType::inu"); + memoryKK->create_kokkos( d_kineticsData.isIntegral, h_kineticsData.isIntegral, nreactions, "KineticsType::isIntegral"); for (int i = 0; i < nreactions; ++i) { @@ -1286,9 +1289,9 @@ void FixRxKokkos<DeviceType>::create_kinetics_data(void) //{ // Dense option - memory->create_kokkos( d_kineticsData.stoich, h_kineticsData.stoich, nreactions, nspecies, "KineticsType::stoich"); - memory->create_kokkos( d_kineticsData.stoichReactants, h_kineticsData.stoichReactants, nreactions, nspecies, "KineticsType::stoichReactants"); - memory->create_kokkos( d_kineticsData.stoichProducts, h_kineticsData.stoichProducts, nreactions, nspecies, "KineticsType::stoichProducts"); + memoryKK->create_kokkos( d_kineticsData.stoich, h_kineticsData.stoich, nreactions, nspecies, "KineticsType::stoich"); + memoryKK->create_kokkos( d_kineticsData.stoichReactants, h_kineticsData.stoichReactants, nreactions, nspecies, "KineticsType::stoichReactants"); + memoryKK->create_kokkos( d_kineticsData.stoichProducts, h_kineticsData.stoichProducts, nreactions, nspecies, "KineticsType::stoichProducts"); for (int i = 0; i < nreactions; ++i) for (int k = 0; k < nspecies; ++k) @@ -1445,8 +1448,8 @@ void FixRxKokkos<DeviceType>::solve_reactions(const int vflag, const bool isPreF const int count = nlocal + (newton_pair ? nghost : 0); if (count > k_dpdThetaLocal.template view<DeviceType>().dimension_0()) { - memory->destroy_kokkos (k_dpdThetaLocal, dpdThetaLocal); - memory->create_kokkos (k_dpdThetaLocal, dpdThetaLocal, count, "FixRxKokkos::dpdThetaLocal"); + memoryKK->destroy_kokkos (k_dpdThetaLocal, dpdThetaLocal); + memoryKK->create_kokkos (k_dpdThetaLocal, dpdThetaLocal, count, "FixRxKokkos::dpdThetaLocal"); this->d_dpdThetaLocal = k_dpdThetaLocal.template view<DeviceType>(); this->h_dpdThetaLocal = k_dpdThetaLocal.h_view; } @@ -1511,8 +1514,8 @@ void FixRxKokkos<DeviceType>::solve_reactions(const int vflag, const bool isPreF if (odeIntegrationFlag == ODE_LAMMPS_RKF45 && diagnosticFrequency == 1) { - memory->create_kokkos (k_diagnosticCounterPerODEnSteps, diagnosticCounterPerODEnSteps, nlocal, "FixRxKokkos::diagnosticCounterPerODEnSteps"); - memory->create_kokkos (k_diagnosticCounterPerODEnFuncs, diagnosticCounterPerODEnFuncs, nlocal, "FixRxKokkos::diagnosticCounterPerODEnFuncs"); + memoryKK->create_kokkos (k_diagnosticCounterPerODEnSteps, diagnosticCounterPerODEnSteps, nlocal, "FixRxKokkos::diagnosticCounterPerODEnSteps"); + memoryKK->create_kokkos (k_diagnosticCounterPerODEnFuncs, diagnosticCounterPerODEnFuncs, nlocal, "FixRxKokkos::diagnosticCounterPerODEnFuncs"); d_diagnosticCounterPerODEnSteps = k_diagnosticCounterPerODEnSteps.template view<DeviceType>(); d_diagnosticCounterPerODEnFuncs = k_diagnosticCounterPerODEnFuncs.template view<DeviceType>(); @@ -1542,8 +1545,8 @@ void FixRxKokkos<DeviceType>::solve_reactions(const int vflag, const bool isPreF //typename ArrayTypes<DeviceType>::t_double_1d d_scratchSpace("d_scratchSpace", scratchSpaceSize * nlocal); if (nlocal*scratchSpaceSize > d_scratchSpace.dimension_0()) { - memory->destroy_kokkos (d_scratchSpace); - memory->create_kokkos (d_scratchSpace, nlocal*scratchSpaceSize, "FixRxKokkos::d_scratchSpace"); + memoryKK->destroy_kokkos (d_scratchSpace); + memoryKK->create_kokkos (d_scratchSpace, nlocal*scratchSpaceSize, "FixRxKokkos::d_scratchSpace"); } #if 0 @@ -1811,8 +1814,8 @@ void FixRxKokkos<DeviceType>::odeDiagnostics(void) my_min[FuncSum] = std::min( my_min[FuncSum], (double)nFuncs ); } - memory->destroy_kokkos( k_diagnosticCounterPerODEnSteps, diagnosticCounterPerODEnSteps ); - memory->destroy_kokkos( k_diagnosticCounterPerODEnFuncs, diagnosticCounterPerODEnFuncs ); + memoryKK->destroy_kokkos( k_diagnosticCounterPerODEnSteps, diagnosticCounterPerODEnSteps ); + memoryKK->destroy_kokkos( k_diagnosticCounterPerODEnFuncs, diagnosticCounterPerODEnFuncs ); MPI_Reduce (my_sum_sq, sum_sq, 2*numCounters, MPI_DOUBLE, MPI_SUM, 0, world); @@ -2022,10 +2025,10 @@ void FixRxKokkos<DeviceType>::computeLocalTemperature() { const int ntypes = atom->ntypes; - //memory->create_kokkos (k_cutsq, h_cutsq, ntypes+1, ntypes+1, "pair:cutsq"); + //memoryKK->create_kokkos (k_cutsq, h_cutsq, ntypes+1, ntypes+1, "pair:cutsq"); if (ntypes+1 > k_cutsq.dimension_0()) { - memory->destroy_kokkos (k_cutsq); - memory->create_kokkos (k_cutsq, ntypes+1, ntypes+1, "FixRxKokkos::k_cutsq"); + memoryKK->destroy_kokkos (k_cutsq); + memoryKK->create_kokkos (k_cutsq, ntypes+1, ntypes+1, "FixRxKokkos::k_cutsq"); d_cutsq = k_cutsq.template view<DeviceType>(); } @@ -2043,10 +2046,10 @@ void FixRxKokkos<DeviceType>::computeLocalTemperature() // Initialize the local temperature weight array int sumWeightsCt = nlocal + (NEWTON_PAIR ? nghost : 0); - //memory->create_kokkos (k_sumWeights, sumWeights, sumWeightsCt, "FixRxKokkos::sumWeights"); + //memoryKK->create_kokkos (k_sumWeights, sumWeights, sumWeightsCt, "FixRxKokkos::sumWeights"); if (sumWeightsCt > k_sumWeights.template view<DeviceType>().dimension_0()) { - memory->destroy_kokkos(k_sumWeights, sumWeights); - memory->create_kokkos (k_sumWeights, sumWeightsCt, "FixRxKokkos::sumWeights"); + memoryKK->destroy_kokkos(k_sumWeights, sumWeights); + memoryKK->create_kokkos (k_sumWeights, sumWeightsCt, "FixRxKokkos::sumWeights"); d_sumWeights = k_sumWeights.template view<DeviceType>(); h_sumWeights = k_sumWeights.h_view; } diff --git a/src/KOKKOS/fix_setforce_kokkos.cpp b/src/KOKKOS/fix_setforce_kokkos.cpp index 2ef04ad6ee3285534ab1f36a26842641fcf5802c..e99160989497381ef63985fad512ec3e2d3af664 100644 --- a/src/KOKKOS/fix_setforce_kokkos.cpp +++ b/src/KOKKOS/fix_setforce_kokkos.cpp @@ -22,10 +22,11 @@ #include "respa.h" #include "input.h" #include "variable.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "force.h" #include "atom_masks.h" +#include "kokkos_base.h" using namespace LAMMPS_NS; using namespace FixConst; @@ -45,7 +46,7 @@ FixSetForceKokkos<DeviceType>::FixSetForceKokkos(LAMMPS *lmp, int narg, char **a datamask_modify = EMPTY_MASK; memory->destroy(sforce); - memory->create_kokkos(k_sforce,sforce,maxatom,3,"setforce:sforce"); + memoryKK->create_kokkos(k_sforce,sforce,maxatom,3,"setforce:sforce"); } /* ---------------------------------------------------------------------- */ @@ -55,7 +56,7 @@ FixSetForceKokkos<DeviceType>::~FixSetForceKokkos() { if (copymode) return; - memory->destroy_kokkos(k_sforce,sforce); + memoryKK->destroy_kokkos(k_sforce,sforce); sforce = NULL; } @@ -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 = dynamic_cast<KokkosBase*>(region); + regionKKBase->match_all_kokkos(groupbit,k_match); k_match.template sync<DeviceType>(); d_match = k_match.template view<DeviceType>(); } @@ -99,8 +101,8 @@ void FixSetForceKokkos<DeviceType>::post_force(int vflag) if (varflag == ATOM && atom->nmax > maxatom) { maxatom = atom->nmax; - memory->destroy_kokkos(k_sforce,sforce); - memory->create_kokkos(k_sforce,sforce,maxatom,3,"setforce:sforce"); + memoryKK->destroy_kokkos(k_sforce,sforce); + memoryKK->create_kokkos(k_sforce,sforce,maxatom,3,"setforce:sforce"); } foriginal[0] = foriginal[1] = foriginal[2] = 0.0; diff --git a/src/KOKKOS/fix_shardlow_kokkos.cpp b/src/KOKKOS/fix_shardlow_kokkos.cpp index 98bbb02714ca454b01e4ac3b65351b3d029ac06e..cc1bd6bede9fabe94f270464cd1ead9558ecad8e 100644 --- a/src/KOKKOS/fix_shardlow_kokkos.cpp +++ b/src/KOKKOS/fix_shardlow_kokkos.cpp @@ -50,7 +50,7 @@ #include "neighbor.h" #include "neigh_list_kokkos.h" #include "neigh_request.h" -#include "memory.h" +#include "memory_kokkos.h" #include "domain.h" #include "modify.h" // #include "pair_dpd_fdt.h" diff --git a/src/KOKKOS/gridcomm_kokkos.cpp b/src/KOKKOS/gridcomm_kokkos.cpp index 6871ef67aee6464b442e67e21dc0776e486d5a43..fdfaf296ef335f30654c1aea300db1102dfca772 100644 --- a/src/KOKKOS/gridcomm_kokkos.cpp +++ b/src/KOKKOS/gridcomm_kokkos.cpp @@ -15,8 +15,9 @@ #include "gridcomm_kokkos.h" #include "comm.h" #include "kspace.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" +#include "kokkos_base.h" using namespace LAMMPS_NS; @@ -126,8 +127,8 @@ template<class DeviceType> GridCommKokkos<DeviceType>::~GridCommKokkos() { for (int i = 0; i < nswap; i++) { - //memory->destroy_kokkos(swap[i].k_packlist,swap[i].packlist); - //memory->destroy_kokkos(swap[i].k_unpacklist,swap[i].unpacklist); + //memoryKK->destroy_kokkos(swap[i].k_packlist,swap[i].packlist); + //memoryKK->destroy_kokkos(swap[i].k_unpacklist,swap[i].unpacklist); } memory->sfree(swap); @@ -515,11 +516,13 @@ void GridCommKokkos<DeviceType>::forward_comm(KSpace *kspace, int which) k_packlist.sync<DeviceType>(); k_unpacklist.sync<DeviceType>(); + KokkosBase* kspaceKKBase = dynamic_cast<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 = dynamic_cast<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/improper_class2_kokkos.cpp b/src/KOKKOS/improper_class2_kokkos.cpp index c2cb7dfe2bf55df1f64988abe60a152d42984654..d2d465a250f1ff97bfd4668fe26eec489f62f850 100644 --- a/src/KOKKOS/improper_class2_kokkos.cpp +++ b/src/KOKKOS/improper_class2_kokkos.cpp @@ -26,7 +26,7 @@ #include "force.h" #include "update.h" #include "math_const.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -58,8 +58,8 @@ template<class DeviceType> ImproperClass2Kokkos<DeviceType>::~ImproperClass2Kokkos() { if (!copymode) { - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); } } @@ -78,15 +78,15 @@ void ImproperClass2Kokkos<DeviceType>::compute(int eflag_in, int vflag_in) if (eflag_atom) { //if(k_eatom.dimension_0()<maxeatom) { // won't work without adding zero functor - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom"); d_eatom = k_eatom.template view<DeviceType>(); //} } if (vflag_atom) { //if(k_vatom.dimension_0()<maxvatom) { // won't work without adding zero functor - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"improper:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"improper:vatom"); d_vatom = k_vatom.template view<DeviceType>(); //} } diff --git a/src/KOKKOS/improper_harmonic_kokkos.cpp b/src/KOKKOS/improper_harmonic_kokkos.cpp index 73e105864f630b0516ac90ba41058cf101878bea..49dd36ed194da6ea94df99144f75602ed11af5ea 100644 --- a/src/KOKKOS/improper_harmonic_kokkos.cpp +++ b/src/KOKKOS/improper_harmonic_kokkos.cpp @@ -26,7 +26,7 @@ #include "force.h" #include "update.h" #include "math_const.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -58,8 +58,8 @@ template<class DeviceType> ImproperHarmonicKokkos<DeviceType>::~ImproperHarmonicKokkos() { if (!copymode) { - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); } } @@ -78,15 +78,15 @@ void ImproperHarmonicKokkos<DeviceType>::compute(int eflag_in, int vflag_in) if (eflag_atom) { //if(k_eatom.dimension_0()<maxeatom) { // won't work without adding zero functor - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"improper:eatom"); d_eatom = k_eatom.template view<DeviceType>(); //} } if (vflag_atom) { //if(k_vatom.dimension_0()<maxvatom) { // won't work without adding zero functor - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"improper:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"improper:vatom"); d_vatom = k_vatom.template view<DeviceType>(); //} } diff --git a/src/KOKKOS/kokkos.cpp b/src/KOKKOS/kokkos.cpp index 2b02624dcef30f29947183cb4f2a5acb68916779..5a742233534403a48b99ca12eae8f8873583401f 100644 --- a/src/KOKKOS/kokkos.cpp +++ b/src/KOKKOS/kokkos.cpp @@ -23,6 +23,7 @@ #include "neighbor_kokkos.h" #include "neigh_list_kokkos.h" #include "error.h" +#include "memory_kokkos.h" using namespace LAMMPS_NS; @@ -33,6 +34,10 @@ KokkosLMP::KokkosLMP(LAMMPS *lmp, int narg, char **arg) : Pointers(lmp) kokkos_exists = 1; lmp->kokkos = this; + delete memory; + memory = new MemoryKokkos(lmp); + memoryKK = (MemoryKokkos*) memory; + auto_sync = 1; int me = 0; diff --git a/src/KOKKOS/kokkos_base.h b/src/KOKKOS/kokkos_base.h new file mode 100644 index 0000000000000000000000000000000000000000..3279cb2947272022704036cbd51ddcb10ff788ef --- /dev/null +++ b/src/KOKKOS/kokkos_base.h @@ -0,0 +1,47 @@ +/* ---------------------------------------------------------------------- + LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator + http://lammps.sandia.gov, Sandia National Laboratories + Steve Plimpton, sjplimp@sandia.gov + + Copyright (2003) Sandia Corporation. Under the terms of Contract + DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains + certain rights in this software. This software is distributed under + the GNU General Public License. + + See the README file in the top-level LAMMPS directory. +------------------------------------------------------------------------- */ + +#ifndef KOKKOS_BASE_H +#define KOKKOS_BASE_H + +#include "kokkos_type.h" + +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) {} +}; + +} + +#endif + +/* ERROR/WARNING messages: + +*/ diff --git a/src/KOKKOS/memory_kokkos.h b/src/KOKKOS/memory_kokkos.h index 8ade198c4056727f78b620300b72783f15375431..9f930faae2e4f973a0849064a585888e29a8483f 100644 --- a/src/KOKKOS/memory_kokkos.h +++ b/src/KOKKOS/memory_kokkos.h @@ -11,6 +11,18 @@ See the README file in the top-level LAMMPS directory. ------------------------------------------------------------------------- */ +#ifndef LMP_MEMORY_KOKKOS_H +#define LMP_MEMORY_KOKKOS_H + +#include "memory.h" +#include "kokkos_type.h" + +namespace LAMMPS_NS { + +class MemoryKokkos : public Memory { + public: + MemoryKokkos(class LAMMPS *lmp) : Memory(lmp) {} + /* ---------------------------------------------------------------------- Kokkos versions of create/grow/destroy multi-dimensional arrays ------------------------------------------------------------------------- */ @@ -279,3 +291,10 @@ void destroy_kokkos(TYPE data, typename TYPE::value_type** &array) sfree(array); array = NULL; } + +}; + +} + +#endif + diff --git a/src/KOKKOS/nbin_ssa_kokkos.cpp b/src/KOKKOS/nbin_ssa_kokkos.cpp index ab97cb584805b6328df67cfe954d02ba44fc3907..1a881658567524b181bbeb52af3b97a17246fc12 100644 --- a/src/KOKKOS/nbin_ssa_kokkos.cpp +++ b/src/KOKKOS/nbin_ssa_kokkos.cpp @@ -26,7 +26,7 @@ #include "error.h" #include "atom_masks.h" -// #include "memory.h" +// #include "memory_kokkos.h" using namespace LAMMPS_NS; diff --git a/src/KOKKOS/neigh_bond_kokkos.cpp b/src/KOKKOS/neigh_bond_kokkos.cpp index a674e7cec48a0e16ff9aa91f36ec85c53b23c0a5..3ecc8b5e516c8243a36a668fcdceb7b434e7145c 100644 --- a/src/KOKKOS/neigh_bond_kokkos.cpp +++ b/src/KOKKOS/neigh_bond_kokkos.cpp @@ -24,7 +24,7 @@ #include "domain_kokkos.h" #include "output.h" #include "thermo.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "modify.h" #include "fix.h" @@ -80,27 +80,27 @@ void NeighBondKokkos<DeviceType>::init_topology_kk() { if (atom->molecular && atom->nbonds && maxbond == 0) { if (nprocs == 1) maxbond = atom->nbonds; else maxbond = static_cast<int> (LB_FACTOR * atom->nbonds / nprocs); - memory->create_kokkos(k_bondlist,neighbor->bondlist,maxbond,3,"neigh:neighbor->bondlist"); + memoryKK->create_kokkos(k_bondlist,neighbor->bondlist,maxbond,3,"neigh:neighbor->bondlist"); } if (atom->molecular && atom->nangles && maxangle == 0) { if (nprocs == 1) maxangle = atom->nangles; else maxangle = static_cast<int> (LB_FACTOR * atom->nangles / nprocs); - memory->create_kokkos(k_anglelist,neighbor->anglelist,maxangle,4,"neigh:neighbor->anglelist"); + memoryKK->create_kokkos(k_anglelist,neighbor->anglelist,maxangle,4,"neigh:neighbor->anglelist"); } if (atom->molecular && atom->ndihedrals && maxdihedral == 0) { if (nprocs == 1) maxdihedral = atom->ndihedrals; else maxdihedral = static_cast<int> (LB_FACTOR * atom->ndihedrals / nprocs); - memory->create_kokkos(k_dihedrallist,neighbor->dihedrallist,maxdihedral,5,"neigh:neighbor->dihedrallist"); + memoryKK->create_kokkos(k_dihedrallist,neighbor->dihedrallist,maxdihedral,5,"neigh:neighbor->dihedrallist"); } if (atom->molecular && atom->nimpropers && maximproper == 0) { if (nprocs == 1) maximproper = atom->nimpropers; else maximproper = static_cast<int> (LB_FACTOR * atom->nimpropers / nprocs); - memory->create_kokkos(k_improperlist,neighbor->improperlist,maximproper,5,"neigh:neighbor->improperlist"); + memoryKK->create_kokkos(k_improperlist,neighbor->improperlist,maximproper,5,"neigh:neighbor->improperlist"); } // set flags that determine which topology neighboring routines to use @@ -283,7 +283,7 @@ void NeighBondKokkos<DeviceType>::bond_all() k_fail_flag.template sync<LMPHostType>(); if (h_fail_flag()) { maxbond = neighbor->nbondlist + BONDDELTA; - memory->grow_kokkos(k_bondlist,neighbor->bondlist,maxbond,3,"neighbor:neighbor->bondlist"); + memoryKK->grow_kokkos(k_bondlist,neighbor->bondlist,maxbond,3,"neighbor:neighbor->bondlist"); v_bondlist = k_bondlist.view<DeviceType>(); } } while (h_fail_flag()); @@ -378,7 +378,7 @@ void NeighBondKokkos<DeviceType>::bond_partial() k_fail_flag.template sync<LMPHostType>(); if (h_fail_flag()) { maxbond = neighbor->nbondlist + BONDDELTA; - memory->grow_kokkos(k_bondlist,neighbor->bondlist,maxbond,3,"neighbor:neighbor->bondlist"); + memoryKK->grow_kokkos(k_bondlist,neighbor->bondlist,maxbond,3,"neighbor:neighbor->bondlist"); v_bondlist = k_bondlist.view<DeviceType>(); } } while (h_fail_flag()); @@ -500,7 +500,7 @@ void NeighBondKokkos<DeviceType>::angle_all() k_fail_flag.template sync<LMPHostType>(); if (h_fail_flag()) { maxangle = neighbor->nanglelist + BONDDELTA; - memory->grow_kokkos(k_anglelist,neighbor->anglelist,maxangle,4,"neighbor:neighbor->anglelist"); + memoryKK->grow_kokkos(k_anglelist,neighbor->anglelist,maxangle,4,"neighbor:neighbor->anglelist"); v_anglelist = k_anglelist.view<DeviceType>(); } } while (h_fail_flag()); @@ -602,7 +602,7 @@ void NeighBondKokkos<DeviceType>::angle_partial() k_fail_flag.template sync<LMPHostType>(); if (h_fail_flag()) { maxangle = neighbor->nanglelist + BONDDELTA; - memory->grow_kokkos(k_anglelist,neighbor->anglelist,maxangle,4,"neighbor:neighbor->anglelist"); + memoryKK->grow_kokkos(k_anglelist,neighbor->anglelist,maxangle,4,"neighbor:neighbor->anglelist"); v_anglelist = k_anglelist.view<DeviceType>(); } } while (h_fail_flag()); @@ -744,7 +744,7 @@ void NeighBondKokkos<DeviceType>::dihedral_all() k_fail_flag.template sync<LMPHostType>(); if (h_fail_flag()) { maxdihedral = neighbor->ndihedrallist + BONDDELTA; - memory->grow_kokkos(k_dihedrallist,neighbor->dihedrallist,maxdihedral,5,"neighbor:neighbor->dihedrallist"); + memoryKK->grow_kokkos(k_dihedrallist,neighbor->dihedrallist,maxdihedral,5,"neighbor:neighbor->dihedrallist"); v_dihedrallist = k_dihedrallist.view<DeviceType>(); } } while (h_fail_flag()); @@ -851,7 +851,7 @@ void NeighBondKokkos<DeviceType>::dihedral_partial() k_fail_flag.template sync<LMPHostType>(); if (h_fail_flag()) { maxdihedral = neighbor->ndihedrallist + BONDDELTA; - memory->grow_kokkos(k_dihedrallist,neighbor->dihedrallist,maxdihedral,5,"neighbor:neighbor->dihedrallist"); + memoryKK->grow_kokkos(k_dihedrallist,neighbor->dihedrallist,maxdihedral,5,"neighbor:neighbor->dihedrallist"); v_dihedrallist = k_dihedrallist.view<DeviceType>(); } } while (h_fail_flag()); @@ -1015,7 +1015,7 @@ void NeighBondKokkos<DeviceType>::improper_all() k_fail_flag.template sync<LMPHostType>(); if (h_fail_flag()) { maximproper = neighbor->nimproperlist + BONDDELTA; - memory->grow_kokkos(k_improperlist,neighbor->improperlist,maximproper,5,"neighbor:neighbor->improperlist"); + memoryKK->grow_kokkos(k_improperlist,neighbor->improperlist,maximproper,5,"neighbor:neighbor->improperlist"); v_improperlist = k_improperlist.view<DeviceType>(); } } while (h_fail_flag()); @@ -1122,7 +1122,7 @@ void NeighBondKokkos<DeviceType>::improper_partial() k_fail_flag.template sync<LMPHostType>(); if (h_fail_flag()) { maximproper = neighbor->nimproperlist + BONDDELTA; - memory->grow_kokkos(k_improperlist,neighbor->improperlist,maximproper,5,"neighbor:neighbor->improperlist"); + memoryKK->grow_kokkos(k_improperlist,neighbor->improperlist,maximproper,5,"neighbor:neighbor->improperlist"); v_improperlist = k_improperlist.view<DeviceType>(); } } while (h_fail_flag()); diff --git a/src/KOKKOS/neigh_list_kokkos.cpp b/src/KOKKOS/neigh_list_kokkos.cpp index 04454e53cb302e2ad696809c034e1925a8aaa83d..98294a802aaa88e16a2c1ca8269cf235d149ff3a 100644 --- a/src/KOKKOS/neigh_list_kokkos.cpp +++ b/src/KOKKOS/neigh_list_kokkos.cpp @@ -13,7 +13,7 @@ #include "neigh_list_kokkos.h" #include "atom.h" -#include "memory.h" +#include "memory_kokkos.h" using namespace LAMMPS_NS; diff --git a/src/KOKKOS/neighbor_kokkos.cpp b/src/KOKKOS/neighbor_kokkos.cpp index f34b149864d1c24179fdf2f41be36143ed5f9807..8d36add10bd6a10d8d861cd0cff5a27cef5fd1f3 100644 --- a/src/KOKKOS/neighbor_kokkos.cpp +++ b/src/KOKKOS/neighbor_kokkos.cpp @@ -12,11 +12,11 @@ ------------------------------------------------------------------------- */ #include "neighbor_kokkos.h" -#include "atom.h" +#include "atom_kokkos.h" #include "pair.h" #include "fix.h" #include "neigh_request.h" -#include "memory.h" +#include "memory_kokkos.h" #include "update.h" #include "atom_masks.h" #include "error.h" @@ -52,24 +52,24 @@ NeighborKokkos::NeighborKokkos(LAMMPS *lmp) : Neighbor(lmp), NeighborKokkos::~NeighborKokkos() { if (!copymode) { - memory->destroy_kokkos(k_cutneighsq,cutneighsq); + memoryKK->destroy_kokkos(k_cutneighsq,cutneighsq); cutneighsq = NULL; - memory->destroy_kokkos(k_ex_type,ex_type); - memory->destroy_kokkos(k_ex1_type,ex1_type); - memory->destroy_kokkos(k_ex2_type,ex2_type); - memory->destroy_kokkos(k_ex1_group,ex1_group); - memory->destroy_kokkos(k_ex2_group,ex2_group); - memory->destroy_kokkos(k_ex_mol_group,ex_mol_group); - memory->destroy_kokkos(k_ex1_bit,ex1_bit); - memory->destroy_kokkos(k_ex2_bit,ex2_bit); - memory->destroy_kokkos(k_ex_mol_bit,ex_mol_bit); - memory->destroy_kokkos(k_ex_mol_intra,ex_mol_intra); - - memory->destroy_kokkos(k_bondlist,bondlist); - memory->destroy_kokkos(k_anglelist,anglelist); - memory->destroy_kokkos(k_dihedrallist,dihedrallist); - memory->destroy_kokkos(k_improperlist,improperlist); + memoryKK->destroy_kokkos(k_ex_type,ex_type); + memoryKK->destroy_kokkos(k_ex1_type,ex1_type); + memoryKK->destroy_kokkos(k_ex2_type,ex2_type); + memoryKK->destroy_kokkos(k_ex1_group,ex1_group); + memoryKK->destroy_kokkos(k_ex2_group,ex2_group); + memoryKK->destroy_kokkos(k_ex_mol_group,ex_mol_group); + memoryKK->destroy_kokkos(k_ex1_bit,ex1_bit); + memoryKK->destroy_kokkos(k_ex2_bit,ex2_bit); + memoryKK->destroy_kokkos(k_ex_mol_bit,ex_mol_bit); + memoryKK->destroy_kokkos(k_ex_mol_intra,ex_mol_intra); + + memoryKK->destroy_kokkos(k_bondlist,bondlist); + memoryKK->destroy_kokkos(k_anglelist,anglelist); + memoryKK->destroy_kokkos(k_dihedrallist,dihedrallist); + memoryKK->destroy_kokkos(k_improperlist,improperlist); } } @@ -90,7 +90,7 @@ void NeighborKokkos::init() void NeighborKokkos::init_cutneighsq_kokkos(int n) { - memory->create_kokkos(k_cutneighsq,cutneighsq,n+1,n+1,"neigh:cutneighsq"); + memoryKK->create_kokkos(k_cutneighsq,cutneighsq,n+1,n+1,"neigh:cutneighsq"); k_cutneighsq.modify<LMPHostType>(); } @@ -112,7 +112,7 @@ void NeighborKokkos::create_kokkos_list(int i) void NeighborKokkos::init_ex_type_kokkos(int n) { - memory->create_kokkos(k_ex_type,ex_type,n+1,n+1,"neigh:ex_type"); + memoryKK->create_kokkos(k_ex_type,ex_type,n+1,n+1,"neigh:ex_type"); k_ex_type.modify<LMPHostType>(); } @@ -120,9 +120,9 @@ void NeighborKokkos::init_ex_type_kokkos(int n) void NeighborKokkos::init_ex_bit_kokkos() { - memory->create_kokkos(k_ex1_bit, ex1_bit, nex_group, "neigh:ex1_bit"); + memoryKK->create_kokkos(k_ex1_bit, ex1_bit, nex_group, "neigh:ex1_bit"); k_ex1_bit.modify<LMPHostType>(); - memory->create_kokkos(k_ex2_bit, ex2_bit, nex_group, "neigh:ex2_bit"); + memoryKK->create_kokkos(k_ex2_bit, ex2_bit, nex_group, "neigh:ex2_bit"); k_ex2_bit.modify<LMPHostType>(); } @@ -130,7 +130,7 @@ void NeighborKokkos::init_ex_bit_kokkos() void NeighborKokkos::init_ex_mol_bit_kokkos() { - memory->create_kokkos(k_ex_mol_bit, ex_mol_bit, nex_mol, "neigh:ex_mol_bit"); + memoryKK->create_kokkos(k_ex_mol_bit, ex_mol_bit, nex_mol, "neigh:ex_mol_bit"); k_ex_mol_bit.modify<LMPHostType>(); } @@ -138,7 +138,7 @@ void NeighborKokkos::init_ex_mol_bit_kokkos() void NeighborKokkos::grow_ex_mol_intra_kokkos() { - memory->grow_kokkos(k_ex_mol_intra, ex_mol_intra, maxex_mol, "neigh:ex_mol_intra"); + memoryKK->grow_kokkos(k_ex_mol_intra, ex_mol_intra, maxex_mol, "neigh:ex_mol_intra"); k_ex_mol_intra.modify<LMPHostType>(); } @@ -335,29 +335,29 @@ void NeighborKokkos::operator()(TagNeighborXhold<DeviceType>, const int &i) cons /* ---------------------------------------------------------------------- */ void NeighborKokkos::modify_ex_type_grow_kokkos(){ - memory->grow_kokkos(k_ex1_type,ex1_type,maxex_type,"neigh:ex1_type"); + memoryKK->grow_kokkos(k_ex1_type,ex1_type,maxex_type,"neigh:ex1_type"); k_ex1_type.modify<LMPHostType>(); - memory->grow_kokkos(k_ex2_type,ex2_type,maxex_type,"neigh:ex2_type"); + memoryKK->grow_kokkos(k_ex2_type,ex2_type,maxex_type,"neigh:ex2_type"); k_ex2_type.modify<LMPHostType>(); } /* ---------------------------------------------------------------------- */ void NeighborKokkos::modify_ex_group_grow_kokkos(){ - memory->grow_kokkos(k_ex1_group,ex1_group,maxex_group,"neigh:ex1_group"); + memoryKK->grow_kokkos(k_ex1_group,ex1_group,maxex_group,"neigh:ex1_group"); k_ex1_group.modify<LMPHostType>(); - memory->grow_kokkos(k_ex2_group,ex2_group,maxex_group,"neigh:ex2_group"); + memoryKK->grow_kokkos(k_ex2_group,ex2_group,maxex_group,"neigh:ex2_group"); k_ex2_group.modify<LMPHostType>(); } /* ---------------------------------------------------------------------- */ void NeighborKokkos::modify_mol_group_grow_kokkos(){ - memory->grow_kokkos(k_ex_mol_group,ex_mol_group,maxex_mol,"neigh:ex_mol_group"); + memoryKK->grow_kokkos(k_ex_mol_group,ex_mol_group,maxex_mol,"neigh:ex_mol_group"); k_ex_mol_group.modify<LMPHostType>(); } /* ---------------------------------------------------------------------- */ void NeighborKokkos::modify_mol_intra_grow_kokkos(){ - memory->grow_kokkos(k_ex_mol_intra,ex_mol_intra,maxex_mol,"neigh:ex_mol_intra"); + memoryKK->grow_kokkos(k_ex_mol_intra,ex_mol_intra,maxex_mol,"neigh:ex_mol_intra"); k_ex_mol_intra.modify<LMPHostType>(); } diff --git a/src/KOKKOS/pair_buck_coul_cut_kokkos.cpp b/src/KOKKOS/pair_buck_coul_cut_kokkos.cpp index 0da8a0a3d6a9267bc6e45a74aa4b52f2f2471c47..ba3eda64ddee4e389b771f5e47168c98c5a6fa08 100644 --- a/src/KOKKOS/pair_buck_coul_cut_kokkos.cpp +++ b/src/KOKKOS/pair_buck_coul_cut_kokkos.cpp @@ -31,7 +31,7 @@ #include "integrate.h" #include "respa.h" #include "math_const.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -65,8 +65,8 @@ PairBuckCoulCutKokkos<DeviceType>::~PairBuckCoulCutKokkos() { if (!copymode) { - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); k_cutsq = DAT::tdual_ffloat_2d(); k_cut_ljsq = DAT::tdual_ffloat_2d(); k_cut_coulsq = DAT::tdual_ffloat_2d(); @@ -98,13 +98,13 @@ void PairBuckCoulCutKokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.view<DeviceType>(); } @@ -252,13 +252,13 @@ void PairBuckCoulCutKokkos<DeviceType>::allocate() int n = atom->ntypes; memory->destroy(cutsq); - memory->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); + memoryKK->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); d_cutsq = k_cutsq.template view<DeviceType>(); memory->destroy(cut_ljsq); - memory->create_kokkos(k_cut_ljsq,cut_ljsq,n+1,n+1,"pair:cut_ljsq"); + memoryKK->create_kokkos(k_cut_ljsq,cut_ljsq,n+1,n+1,"pair:cut_ljsq"); d_cut_ljsq = k_cut_ljsq.template view<DeviceType>(); memory->destroy(cut_coulsq); - memory->create_kokkos(k_cut_coulsq,cut_coulsq,n+1,n+1,"pair:cut_coulsq"); + memoryKK->create_kokkos(k_cut_coulsq,cut_coulsq,n+1,n+1,"pair:cut_coulsq"); d_cut_coulsq = k_cut_coulsq.template view<DeviceType>(); k_params = Kokkos::DualView<params_buck_coul**,Kokkos::LayoutRight,DeviceType>("PairBuckCoulCut::params",n+1,n+1); params = k_params.template view<DeviceType>(); diff --git a/src/KOKKOS/pair_buck_coul_long_kokkos.cpp b/src/KOKKOS/pair_buck_coul_long_kokkos.cpp index 3a5cbd868f725c54a95c6c6cb0fc9a2335137ea5..19af349a63588a464f4e1785be291510fd91acfa 100644 --- a/src/KOKKOS/pair_buck_coul_long_kokkos.cpp +++ b/src/KOKKOS/pair_buck_coul_long_kokkos.cpp @@ -31,7 +31,7 @@ #include "integrate.h" #include "respa.h" #include "math_const.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -73,8 +73,8 @@ template<class DeviceType> PairBuckCoulLongKokkos<DeviceType>::~PairBuckCoulLongKokkos() { if (!copymode) { - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); k_cutsq = DAT::tdual_ffloat_2d(); k_cut_ljsq = DAT::tdual_ffloat_2d(); k_cut_coulsq = DAT::tdual_ffloat_2d(); @@ -117,13 +117,13 @@ void PairBuckCoulLongKokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.view<DeviceType>(); } @@ -308,14 +308,14 @@ void PairBuckCoulLongKokkos<DeviceType>::allocate() int n = atom->ntypes; memory->destroy(cutsq); - memory->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); + memoryKK->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); d_cutsq = k_cutsq.template view<DeviceType>(); memory->destroy(cut_ljsq); - memory->create_kokkos(k_cut_ljsq,cut_ljsq,n+1,n+1,"pair:cut_ljsq"); + memoryKK->create_kokkos(k_cut_ljsq,cut_ljsq,n+1,n+1,"pair:cut_ljsq"); d_cut_ljsq = k_cut_ljsq.template view<DeviceType>(); - memory->create_kokkos(k_cut_coulsq,n+1,n+1,"pair:cut_coulsq"); + memoryKK->create_kokkos(k_cut_coulsq,n+1,n+1,"pair:cut_coulsq"); d_cut_coulsq = k_cut_coulsq.template view<DeviceType>(); k_params = Kokkos::DualView<params_buck_coul**,Kokkos::LayoutRight,DeviceType>("PairBuckCoulLong::params",n+1,n+1); diff --git a/src/KOKKOS/pair_buck_kokkos.cpp b/src/KOKKOS/pair_buck_kokkos.cpp index e7640471d501bc8be6778c4acb32395e18cf950a..fcf14533dc991d469ad1c0e65951ca80d515b0b2 100644 --- a/src/KOKKOS/pair_buck_kokkos.cpp +++ b/src/KOKKOS/pair_buck_kokkos.cpp @@ -31,7 +31,7 @@ #include "integrate.h" #include "respa.h" #include "math_const.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -61,8 +61,8 @@ template<class DeviceType> PairBuckKokkos<DeviceType>::~PairBuckKokkos() { if (!copymode) { - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); k_cutsq = DAT::tdual_ffloat_2d(); memory->sfree(cutsq); eatom = NULL; @@ -87,13 +87,13 @@ void PairBuckKokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.view<DeviceType>(); } @@ -194,7 +194,7 @@ void PairBuckKokkos<DeviceType>::allocate() int n = atom->ntypes; memory->destroy(cutsq); - memory->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); + memoryKK->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); d_cutsq = k_cutsq.template view<DeviceType>(); k_params = Kokkos::DualView<params_buck**,Kokkos::LayoutRight,DeviceType>("PairBuck::params",n+1,n+1); params = k_params.template view<DeviceType>(); diff --git a/src/KOKKOS/pair_coul_cut_kokkos.cpp b/src/KOKKOS/pair_coul_cut_kokkos.cpp index 8edf093e2e118a4315b899aa6c304a9e78f670f9..e20e243c09aefb7313cbe13ce50a498356ea405d 100644 --- a/src/KOKKOS/pair_coul_cut_kokkos.cpp +++ b/src/KOKKOS/pair_coul_cut_kokkos.cpp @@ -27,7 +27,7 @@ #include "integrate.h" #include "respa.h" #include "math_const.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -55,7 +55,7 @@ template<class DeviceType> PairCoulCutKokkos<DeviceType>::~PairCoulCutKokkos() { if (allocated) - memory->destroy_kokkos(k_cutsq, cutsq); + memoryKK->destroy_kokkos(k_cutsq, cutsq); } /* ---------------------------------------------------------------------- */ @@ -86,13 +86,13 @@ void PairCoulCutKokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.view<DeviceType>(); } @@ -190,7 +190,7 @@ void PairCoulCutKokkos<DeviceType>::allocate() int n = atom->ntypes; memory->destroy(cutsq); - memory->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); + memoryKK->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); d_cutsq = k_cutsq.template view<DeviceType>(); k_cut_ljsq = typename ArrayTypes<DeviceType>::tdual_ffloat_2d("pair:cut_ljsq",n+1,n+1); diff --git a/src/KOKKOS/pair_coul_debye_kokkos.cpp b/src/KOKKOS/pair_coul_debye_kokkos.cpp index c331ab8da8f5881441985a0bfdf69d7f9da8bc8d..4cac18cacfcaf2ba3274a3feeb6e6f3d3a65c4bf 100644 --- a/src/KOKKOS/pair_coul_debye_kokkos.cpp +++ b/src/KOKKOS/pair_coul_debye_kokkos.cpp @@ -31,7 +31,7 @@ #include "integrate.h" #include "respa.h" #include "math_const.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -62,7 +62,7 @@ template<class DeviceType> PairCoulDebyeKokkos<DeviceType>::~PairCoulDebyeKokkos() { if (!copymode) { - memory->destroy_kokkos(k_cutsq, cutsq); + memoryKK->destroy_kokkos(k_cutsq, cutsq); } } @@ -93,13 +93,13 @@ void PairCoulDebyeKokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.view<DeviceType>(); } @@ -218,7 +218,7 @@ void PairCoulDebyeKokkos<DeviceType>::allocate() int n = atom->ntypes; memory->destroy(cutsq); - memory->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); + memoryKK->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); d_cutsq = k_cutsq.template view<DeviceType>(); k_params = Kokkos::DualView<params_coul**,Kokkos::LayoutRight,DeviceType>("PairCoulDebye::params",n+1,n+1); params = k_params.template view<DeviceType>(); diff --git a/src/KOKKOS/pair_coul_dsf_kokkos.cpp b/src/KOKKOS/pair_coul_dsf_kokkos.cpp index e6f5407f2d16990085495915562a4c564861aa81..f77e63bbf0551b0412b21e0b7631c098ff082f78 100644 --- a/src/KOKKOS/pair_coul_dsf_kokkos.cpp +++ b/src/KOKKOS/pair_coul_dsf_kokkos.cpp @@ -27,7 +27,7 @@ #include "neighbor.h" #include "neigh_list_kokkos.h" #include "neigh_request.h" -#include "memory.h" +#include "memory_kokkos.h" #include "update.h" #include "integrate.h" #include "respa.h" @@ -65,8 +65,8 @@ template<class DeviceType> PairCoulDSFKokkos<DeviceType>::~PairCoulDSFKokkos() { if (!copymode) { - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); } } @@ -86,13 +86,13 @@ void PairCoulDSFKokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.view<DeviceType>(); } diff --git a/src/KOKKOS/pair_coul_long_kokkos.cpp b/src/KOKKOS/pair_coul_long_kokkos.cpp index 721e140e3374794dad4016fce79888cb60a07145..f2ade3f367c0e01b8118d454a3d0259618cffddf 100644 --- a/src/KOKKOS/pair_coul_long_kokkos.cpp +++ b/src/KOKKOS/pair_coul_long_kokkos.cpp @@ -31,7 +31,7 @@ #include "integrate.h" #include "respa.h" #include "math_const.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -72,8 +72,8 @@ template<class DeviceType> PairCoulLongKokkos<DeviceType>::~PairCoulLongKokkos() { if (!copymode) { - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); k_cutsq = DAT::tdual_ffloat_2d(); k_cut_coulsq = DAT::tdual_ffloat_2d(); memory->sfree(cutsq); @@ -110,13 +110,13 @@ void PairCoulLongKokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.view<DeviceType>(); } @@ -272,10 +272,10 @@ void PairCoulLongKokkos<DeviceType>::allocate() int n = atom->ntypes; memory->destroy(cutsq); - memory->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); + memoryKK->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); d_cutsq = k_cutsq.template view<DeviceType>(); - memory->create_kokkos(k_cut_coulsq,n+1,n+1,"pair:cut_coulsq"); + memoryKK->create_kokkos(k_cut_coulsq,n+1,n+1,"pair:cut_coulsq"); d_cut_coulsq = k_cut_coulsq.template view<DeviceType>(); k_params = Kokkos::DualView<params_coul**,Kokkos::LayoutRight,DeviceType>("PairCoulLong::params",n+1,n+1); diff --git a/src/KOKKOS/pair_coul_wolf_kokkos.cpp b/src/KOKKOS/pair_coul_wolf_kokkos.cpp index 75177e2d81e7a3a9a70ccc754db14651a8d29e56..0f3e9b94290918dd78cb8940cc762d5e82494977 100644 --- a/src/KOKKOS/pair_coul_wolf_kokkos.cpp +++ b/src/KOKKOS/pair_coul_wolf_kokkos.cpp @@ -31,7 +31,7 @@ #include "integrate.h" #include "respa.h" #include "math_const.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -60,8 +60,8 @@ template<class DeviceType> PairCoulWolfKokkos<DeviceType>::~PairCoulWolfKokkos() { if (!copymode) { - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); } } @@ -81,13 +81,13 @@ void PairCoulWolfKokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.view<DeviceType>(); } diff --git a/src/KOKKOS/pair_dpd_fdt_energy_kokkos.cpp b/src/KOKKOS/pair_dpd_fdt_energy_kokkos.cpp index c559ab412f355c1dc10a51adb4308fbf031003d4..7d71719e0a115d99cf25a3a89a91487f9042edc2 100644 --- a/src/KOKKOS/pair_dpd_fdt_energy_kokkos.cpp +++ b/src/KOKKOS/pair_dpd_fdt_energy_kokkos.cpp @@ -28,11 +28,12 @@ #include "neighbor.h" #include "neigh_list.h" #include "neigh_request.h" -#include "memory.h" +#include "memory_kokkos.h" #include "modify.h" #include "pair_dpd_fdt_energy_kokkos.h" #include "error.h" #include "atom_masks.h" +#include "kokkos.h" using namespace LAMMPS_NS; @@ -62,15 +63,15 @@ PairDPDfdtEnergyKokkos<DeviceType>::~PairDPDfdtEnergyKokkos() { if (copymode) return; - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); if (allocated) { - memory->destroy_kokkos(k_duCond,duCond); - memory->destroy_kokkos(k_duMech,duMech); + memoryKK->destroy_kokkos(k_duCond,duCond); + memoryKK->destroy_kokkos(k_duMech,duMech); } - memory->destroy_kokkos(k_cutsq,cutsq); + memoryKK->destroy_kokkos(k_cutsq,cutsq); #ifdef DPD_USE_RAN_MARS rand_pool.destroy(); @@ -167,13 +168,13 @@ void PairDPDfdtEnergyKokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.template view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.template view<DeviceType>(); } @@ -274,11 +275,11 @@ void PairDPDfdtEnergyKokkos<DeviceType>::compute(int eflag_in, int vflag_in) // Allocate memory for duCond and duMech if (allocated) { - memory->destroy_kokkos(k_duCond,duCond); - memory->destroy_kokkos(k_duMech,duMech); + memoryKK->destroy_kokkos(k_duCond,duCond); + memoryKK->destroy_kokkos(k_duMech,duMech); } - memory->create_kokkos(k_duCond,duCond,nlocal+nghost,"pair:duCond"); - memory->create_kokkos(k_duMech,duMech,nlocal+nghost,"pair:duMech"); + memoryKK->create_kokkos(k_duCond,duCond,nlocal+nghost,"pair:duCond"); + memoryKK->create_kokkos(k_duMech,duMech,nlocal+nghost,"pair:duMech"); d_duCond = k_duCond.view<DeviceType>(); d_duMech = k_duMech.view<DeviceType>(); h_duCond = k_duCond.h_view; @@ -641,7 +642,7 @@ void PairDPDfdtEnergyKokkos<DeviceType>::allocate() int nghost = atom->nghost; memory->destroy(cutsq); - memory->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); + memoryKK->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); d_cutsq = k_cutsq.template view<DeviceType>(); k_params = Kokkos::DualView<params_dpd**,Kokkos::LayoutRight,DeviceType>("PairDPDfdtEnergy::params",n+1,n+1); @@ -650,8 +651,8 @@ void PairDPDfdtEnergyKokkos<DeviceType>::allocate() if (!splitFDT_flag) { memory->destroy(duCond); memory->destroy(duMech); - memory->create_kokkos(k_duCond,duCond,nlocal+nghost+1,"pair:duCond"); - memory->create_kokkos(k_duMech,duMech,nlocal+nghost+1,"pair:duMech"); + memoryKK->create_kokkos(k_duCond,duCond,nlocal+nghost+1,"pair:duCond"); + memoryKK->create_kokkos(k_duMech,duMech,nlocal+nghost+1,"pair:duMech"); d_duCond = k_duCond.view<DeviceType>(); d_duMech = k_duMech.view<DeviceType>(); h_duCond = k_duCond.h_view; diff --git a/src/KOKKOS/pair_eam_alloy_kokkos.cpp b/src/KOKKOS/pair_eam_alloy_kokkos.cpp index aa68d0a05438ad2db1256e62cbb979861be47c47..f263cbb6d90be1cd34d5f0f9abab2b4e04174794 100644 --- a/src/KOKKOS/pair_eam_alloy_kokkos.cpp +++ b/src/KOKKOS/pair_eam_alloy_kokkos.cpp @@ -28,7 +28,7 @@ #include "neighbor.h" #include "neigh_list_kokkos.h" #include "neigh_request.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -59,8 +59,8 @@ template<class DeviceType> PairEAMAlloyKokkos<DeviceType>::~PairEAMAlloyKokkos() { if (!copymode) { - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); } } @@ -80,13 +80,13 @@ void PairEAMAlloyKokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.view<DeviceType>(); } diff --git a/src/KOKKOS/pair_eam_alloy_kokkos.h b/src/KOKKOS/pair_eam_alloy_kokkos.h index fb07eec32b1253c61d309f60f1010fb31d4441a6..a962f559085d767f10c1b347a83b9d006c68a548 100644 --- a/src/KOKKOS/pair_eam_alloy_kokkos.h +++ b/src/KOKKOS/pair_eam_alloy_kokkos.h @@ -24,6 +24,7 @@ PairStyle(eam/alloy/kk/host,PairEAMAlloyKokkos<LMPHostType>) #define LMP_PAIR_EAM_ALLOY_KOKKOS_H #include <stdio.h> +#include "kokkos_base.h" #include "pair_kokkos.h" #include "pair_eam.h" #include "neigh_list_kokkos.h" @@ -49,7 +50,7 @@ struct TagPairEAMAlloyKernelC{}; // Cannot use virtual inheritance on the GPU template<class DeviceType> -class PairEAMAlloyKokkos : public PairEAM { +class PairEAMAlloyKokkos : public PairEAM, public KokkosBase { public: enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF}; enum {COUL_FLAG=0}; @@ -59,7 +60,7 @@ class PairEAMAlloyKokkos : public PairEAM { PairEAMAlloyKokkos(class LAMMPS *); virtual ~PairEAMAlloyKokkos(); - virtual void compute(int, int); + void compute(int, int); void init_style(); void *extract(const char *, int &) { return NULL; } void coeff(int, char **); @@ -107,11 +108,11 @@ class PairEAMAlloyKokkos : public PairEAM { const F_FLOAT &epair, const F_FLOAT &fpair, const F_FLOAT &delx, const F_FLOAT &dely, const F_FLOAT &delz) const; - virtual int pack_forward_comm_kokkos(int, DAT::tdual_int_2d, int, DAT::tdual_xfloat_1d&, - int, int *); - virtual void unpack_forward_comm_kokkos(int, int, DAT::tdual_xfloat_1d&); - virtual int pack_forward_comm(int, int *, double *, int, int *); - virtual void unpack_forward_comm(int, int, double *); + int pack_forward_comm_kokkos(int, DAT::tdual_int_2d, int, DAT::tdual_xfloat_1d&, + int, int *); + void unpack_forward_comm_kokkos(int, int, DAT::tdual_xfloat_1d&); + int pack_forward_comm(int, int *, double *, int, int *); + void unpack_forward_comm(int, int, double *); int pack_reverse_comm(int, int, double *); void unpack_reverse_comm(int, int *, double *); @@ -148,7 +149,7 @@ class PairEAMAlloyKokkos : public PairEAM { t_ffloat_2d_n7_randomread d_rhor_spline; t_ffloat_2d_n7_randomread d_z2r_spline; - virtual void file2array(); + void file2array(); void file2array_alloy(); void array2spline(); void interpolate(int, double, double *, t_host_ffloat_2d_n7, int); diff --git a/src/KOKKOS/pair_eam_fs_kokkos.cpp b/src/KOKKOS/pair_eam_fs_kokkos.cpp index a982f94ec4d607ae2e83d1773ed38a98a082ad09..05ec1a644b829cc587ca30a4f99148d78dbeeb11 100644 --- a/src/KOKKOS/pair_eam_fs_kokkos.cpp +++ b/src/KOKKOS/pair_eam_fs_kokkos.cpp @@ -28,7 +28,7 @@ #include "neighbor.h" #include "neigh_list_kokkos.h" #include "neigh_request.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -59,8 +59,8 @@ template<class DeviceType> PairEAMFSKokkos<DeviceType>::~PairEAMFSKokkos() { if (!copymode) { - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); } } @@ -80,13 +80,13 @@ void PairEAMFSKokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.view<DeviceType>(); } diff --git a/src/KOKKOS/pair_eam_fs_kokkos.h b/src/KOKKOS/pair_eam_fs_kokkos.h index d71ec2b8870f103ae1f839d5a90ff791c367afaf..ec87e44ece248f93f0193684311f3587be8ec5d3 100644 --- a/src/KOKKOS/pair_eam_fs_kokkos.h +++ b/src/KOKKOS/pair_eam_fs_kokkos.h @@ -24,6 +24,7 @@ PairStyle(eam/fs/kk/host,PairEAMFSKokkos<LMPHostType>) #define LMP_PAIR_EAM_FS_KOKKOS_H #include <stdio.h> +#include "kokkos_base.h" #include "pair_kokkos.h" #include "pair_eam.h" #include "neigh_list_kokkos.h" @@ -49,7 +50,7 @@ struct TagPairEAMFSKernelC{}; // Cannot use virtual inheritance on the GPU template<class DeviceType> -class PairEAMFSKokkos : public PairEAM { +class PairEAMFSKokkos : public PairEAM, public KokkosBase { public: enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF}; enum {COUL_FLAG=0}; @@ -59,7 +60,7 @@ class PairEAMFSKokkos : public PairEAM { PairEAMFSKokkos(class LAMMPS *); virtual ~PairEAMFSKokkos(); - virtual void compute(int, int); + void compute(int, int); void init_style(); void *extract(const char *, int &) { return NULL; } void coeff(int, char **); @@ -107,11 +108,11 @@ class PairEAMFSKokkos : public PairEAM { const F_FLOAT &epair, const F_FLOAT &fpair, const F_FLOAT &delx, const F_FLOAT &dely, const F_FLOAT &delz) const; - virtual int pack_forward_comm_kokkos(int, DAT::tdual_int_2d, int, DAT::tdual_xfloat_1d&, - int, int *); - virtual void unpack_forward_comm_kokkos(int, int, DAT::tdual_xfloat_1d&); - virtual int pack_forward_comm(int, int *, double *, int, int *); - virtual void unpack_forward_comm(int, int, double *); + int pack_forward_comm_kokkos(int, DAT::tdual_int_2d, int, DAT::tdual_xfloat_1d&, + int, int *); + void unpack_forward_comm_kokkos(int, int, DAT::tdual_xfloat_1d&); + int pack_forward_comm(int, int *, double *, int, int *); + void unpack_forward_comm(int, int, double *); int pack_reverse_comm(int, int, double *); void unpack_reverse_comm(int, int *, double *); @@ -148,7 +149,7 @@ class PairEAMFSKokkos : public PairEAM { t_ffloat_2d_n7_randomread d_rhor_spline; t_ffloat_2d_n7_randomread d_z2r_spline; - virtual void file2array(); + void file2array(); void file2array_fs(); void array2spline(); void interpolate(int, double, double *, t_host_ffloat_2d_n7, int); diff --git a/src/KOKKOS/pair_eam_kokkos.cpp b/src/KOKKOS/pair_eam_kokkos.cpp index 8ac92a1766c62eedbaf37663d5ba1092bbab4599..0aa43b26b2112f6970bae5562b506b4cabe6718d 100644 --- a/src/KOKKOS/pair_eam_kokkos.cpp +++ b/src/KOKKOS/pair_eam_kokkos.cpp @@ -28,7 +28,7 @@ #include "neighbor.h" #include "neigh_list_kokkos.h" #include "neigh_request.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -54,8 +54,8 @@ template<class DeviceType> PairEAMKokkos<DeviceType>::~PairEAMKokkos() { if (!copymode) { - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); } } @@ -75,13 +75,13 @@ void PairEAMKokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.view<DeviceType>(); } diff --git a/src/KOKKOS/pair_eam_kokkos.h b/src/KOKKOS/pair_eam_kokkos.h index 856cc51f7709f24cec67df5d0e8de930f4796c69..34a4795ec54e97986b9cd39e0e79e7932d5d6bd4 100644 --- a/src/KOKKOS/pair_eam_kokkos.h +++ b/src/KOKKOS/pair_eam_kokkos.h @@ -24,6 +24,7 @@ PairStyle(eam/kk/host,PairEAMKokkos<LMPHostType>) #define LMP_PAIR_EAM_KOKKOS_H #include <stdio.h> +#include "kokkos_base.h" #include "pair_kokkos.h" #include "pair_eam.h" #include "neigh_list_kokkos.h" @@ -47,7 +48,7 @@ template<int NEIGHFLAG, int NEWTON_PAIR, int EVFLAG> struct TagPairEAMKernelC{}; template<class DeviceType> -class PairEAMKokkos : public PairEAM { +class PairEAMKokkos : public PairEAM, public KokkosBase { public: enum {EnabledNeighFlags=FULL|HALFTHREAD|HALF}; enum {COUL_FLAG=0}; @@ -57,7 +58,7 @@ class PairEAMKokkos : public PairEAM { PairEAMKokkos(class LAMMPS *); virtual ~PairEAMKokkos(); - virtual void compute(int, int); + void compute(int, int); void init_style(); void *extract(const char *, int &) { return NULL; } @@ -104,11 +105,11 @@ class PairEAMKokkos : public PairEAM { const F_FLOAT &epair, const F_FLOAT &fpair, const F_FLOAT &delx, const F_FLOAT &dely, const F_FLOAT &delz) const; - virtual int pack_forward_comm_kokkos(int, DAT::tdual_int_2d, int, DAT::tdual_xfloat_1d&, - int, int *); - virtual void unpack_forward_comm_kokkos(int, int, DAT::tdual_xfloat_1d&); - virtual int pack_forward_comm(int, int *, double *, int, int *); - virtual void unpack_forward_comm(int, int, double *); + int pack_forward_comm_kokkos(int, DAT::tdual_int_2d, int, DAT::tdual_xfloat_1d&, + int, int *); + void unpack_forward_comm_kokkos(int, int, DAT::tdual_xfloat_1d&); + int pack_forward_comm(int, int *, double *, int, int *); + void unpack_forward_comm(int, int, double *); int pack_reverse_comm(int, int, double *); void unpack_reverse_comm(int, int *, double *); @@ -146,7 +147,7 @@ class PairEAMKokkos : public PairEAM { t_ffloat_2d_n7_randomread d_z2r_spline; void interpolate(int, double, double *, t_host_ffloat_2d_n7, int); - virtual void file2array(); + void file2array(); void array2spline(); typename AT::t_neighbors_2d d_neighbors; diff --git a/src/KOKKOS/pair_exp6_rx_kokkos.cpp b/src/KOKKOS/pair_exp6_rx_kokkos.cpp index 8d65be23af00e7c37e4ce2ad0f1f781a822fc6c4..a09a1d27e34c544c15c79de45ff28ee6a2650e54 100644 --- a/src/KOKKOS/pair_exp6_rx_kokkos.cpp +++ b/src/KOKKOS/pair_exp6_rx_kokkos.cpp @@ -26,13 +26,15 @@ #include "neigh_list.h" #include "math_const.h" #include "math_special_kokkos.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "modify.h" #include "fix.h" #include <float.h> #include "atom_masks.h" #include "neigh_request.h" +#include "atom_kokkos.h" +#include "kokkos.h" using namespace LAMMPS_NS; using namespace MathConst; @@ -89,18 +91,18 @@ PairExp6rxKokkos<DeviceType>::~PairExp6rxKokkos() { if (copymode) return; - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); - memory->destroy_kokkos(k_cutsq,cutsq); + memoryKK->destroy_kokkos(k_cutsq,cutsq); for (int i=0; i < nparams; ++i) { delete[] params[i].name; delete[] params[i].potential; } - memory->destroy_kokkos(k_params,params); + memoryKK->destroy_kokkos(k_params,params); - memory->destroy_kokkos(k_mol2param,mol2param); + memoryKK->destroy_kokkos(k_mol2param,mol2param); } /* ---------------------------------------------------------------------- */ @@ -151,13 +153,13 @@ void PairExp6rxKokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.template view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.template view<DeviceType>(); } @@ -1660,7 +1662,7 @@ void PairExp6rxKokkos<DeviceType>::allocate() for (int j = i; j <= ntypes; j++) setflag[i][j] = 0; - memory->create_kokkos(k_cutsq,cutsq,ntypes+1,ntypes+1,"pair:cutsq"); + memoryKK->create_kokkos(k_cutsq,cutsq,ntypes+1,ntypes+1,"pair:cutsq"); d_cutsq = k_cutsq.template view<DeviceType>(); k_cutsq.template modify<LMPHostType>(); @@ -1697,7 +1699,7 @@ void PairExp6rxKokkos<DeviceType>::read_file(char *file) int params_per_line = 5; char **words = new char*[params_per_line+1]; - memory->destroy_kokkos(k_params,params); + memoryKK->destroy_kokkos(k_params,params); params = NULL; nparams = maxparam = 0; @@ -1777,7 +1779,7 @@ void PairExp6rxKokkos<DeviceType>::read_file(char *file) if (nparams == maxparam) { k_params.template modify<LMPHostType>(); maxparam += DELTA; - memory->grow_kokkos(k_params,params,maxparam, + memoryKK->grow_kokkos(k_params,params,maxparam, "pair:params"); } @@ -1816,8 +1818,8 @@ void PairExp6rxKokkos<DeviceType>::setup() // set mol2param for all combinations // must be a single exact match to lines read from file - memory->destroy_kokkos(k_mol2param,mol2param); - memory->create_kokkos(k_mol2param,mol2param,nspecies,"pair:mol2param"); + memoryKK->destroy_kokkos(k_mol2param,mol2param); + memoryKK->create_kokkos(k_mol2param,mol2param,nspecies,"pair:mol2param"); for (i = 0; i < nspecies; i++) { n = -1; diff --git a/src/KOKKOS/pair_hybrid_kokkos.cpp b/src/KOKKOS/pair_hybrid_kokkos.cpp index 337b56c6cee7e11e3a03dcd810bd62e161bb6813..db757f6ce1847806d266de13037d99b6489732a3 100644 --- a/src/KOKKOS/pair_hybrid_kokkos.cpp +++ b/src/KOKKOS/pair_hybrid_kokkos.cpp @@ -23,10 +23,11 @@ #include "neigh_request.h" #include "update.h" #include "comm.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "respa.h" #include "atom_masks.h" +#include "kokkos.h" using namespace LAMMPS_NS; diff --git a/src/KOKKOS/pair_kokkos.h b/src/KOKKOS/pair_kokkos.h index f0e357270c1d5b8340c570f4a1d5be4f26d36a27..4c215ed16287d54ed3237aa678612c74c5faabf0 100644 --- a/src/KOKKOS/pair_kokkos.h +++ b/src/KOKKOS/pair_kokkos.h @@ -20,6 +20,7 @@ #include "Kokkos_Macros.hpp" #include "pair.h" +#include "neighbor_kokkos.h" #include "neigh_list_kokkos.h" #include "Kokkos_Vectorization.hpp" diff --git a/src/KOKKOS/pair_lj_charmm_coul_charmm_implicit_kokkos.cpp b/src/KOKKOS/pair_lj_charmm_coul_charmm_implicit_kokkos.cpp index e7cf7ba42ad80ba7166322a2dd1999c3e3c19f3c..7d395cb3b464746f7efeb9e52232f122dcfb8b4d 100644 --- a/src/KOKKOS/pair_lj_charmm_coul_charmm_implicit_kokkos.cpp +++ b/src/KOKKOS/pair_lj_charmm_coul_charmm_implicit_kokkos.cpp @@ -31,7 +31,7 @@ #include "integrate.h" #include "respa.h" #include "math_const.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -73,8 +73,8 @@ template<class DeviceType> PairLJCharmmCoulCharmmImplicitKokkos<DeviceType>::~PairLJCharmmCoulCharmmImplicitKokkos() { if (!copymode) { - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); k_cutsq = DAT::tdual_ffloat_2d(); k_cut_ljsq = DAT::tdual_ffloat_2d(); k_cut_coulsq = DAT::tdual_ffloat_2d(); @@ -118,13 +118,13 @@ void PairLJCharmmCoulCharmmImplicitKokkos<DeviceType>::compute(int eflag_in, int // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.view<DeviceType>(); } @@ -315,14 +315,14 @@ void PairLJCharmmCoulCharmmImplicitKokkos<DeviceType>::allocate() int n = atom->ntypes; memory->destroy(cutsq); - memory->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); + memoryKK->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); d_cutsq = k_cutsq.template view<DeviceType>(); //memory->destroy(cut_ljsq); - memory->create_kokkos(k_cut_ljsq,n+1,n+1,"pair:cut_ljsq"); + memoryKK->create_kokkos(k_cut_ljsq,n+1,n+1,"pair:cut_ljsq"); d_cut_ljsq = k_cut_ljsq.template view<DeviceType>(); - memory->create_kokkos(k_cut_coulsq,n+1,n+1,"pair:cut_coulsq"); + memoryKK->create_kokkos(k_cut_coulsq,n+1,n+1,"pair:cut_coulsq"); d_cut_coulsq = k_cut_coulsq.template view<DeviceType>(); k_params = Kokkos::DualView<params_lj_coul**,Kokkos::LayoutRight,DeviceType>("PairLJCharmmCoulCharmmImplicit::params",n+1,n+1); diff --git a/src/KOKKOS/pair_lj_charmm_coul_charmm_kokkos.cpp b/src/KOKKOS/pair_lj_charmm_coul_charmm_kokkos.cpp index a456d6e27600c36911da43be21f8151a91943148..2663d71a1febe0aced1bacec03b1c80ff6136715 100644 --- a/src/KOKKOS/pair_lj_charmm_coul_charmm_kokkos.cpp +++ b/src/KOKKOS/pair_lj_charmm_coul_charmm_kokkos.cpp @@ -31,7 +31,7 @@ #include "integrate.h" #include "respa.h" #include "math_const.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -73,8 +73,8 @@ template<class DeviceType> PairLJCharmmCoulCharmmKokkos<DeviceType>::~PairLJCharmmCoulCharmmKokkos() { if (!copymode) { - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); k_cutsq = DAT::tdual_ffloat_2d(); k_cut_ljsq = DAT::tdual_ffloat_2d(); k_cut_coulsq = DAT::tdual_ffloat_2d(); @@ -118,13 +118,13 @@ void PairLJCharmmCoulCharmmKokkos<DeviceType>::compute(int eflag_in, int vflag_i // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.view<DeviceType>(); } @@ -316,14 +316,14 @@ void PairLJCharmmCoulCharmmKokkos<DeviceType>::allocate() int n = atom->ntypes; memory->destroy(cutsq); - memory->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); + memoryKK->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); d_cutsq = k_cutsq.template view<DeviceType>(); //memory->destroy(cut_ljsq); - memory->create_kokkos(k_cut_ljsq,n+1,n+1,"pair:cut_ljsq"); + memoryKK->create_kokkos(k_cut_ljsq,n+1,n+1,"pair:cut_ljsq"); d_cut_ljsq = k_cut_ljsq.template view<DeviceType>(); - memory->create_kokkos(k_cut_coulsq,n+1,n+1,"pair:cut_coulsq"); + memoryKK->create_kokkos(k_cut_coulsq,n+1,n+1,"pair:cut_coulsq"); d_cut_coulsq = k_cut_coulsq.template view<DeviceType>(); k_params = Kokkos::DualView<params_lj_coul**,Kokkos::LayoutRight,DeviceType>("PairLJCharmmCoulCharmm::params",n+1,n+1); diff --git a/src/KOKKOS/pair_lj_charmm_coul_long_kokkos.cpp b/src/KOKKOS/pair_lj_charmm_coul_long_kokkos.cpp index dffbbb638ff0d1f25819d99b504e58d2d9eae300..81271c7d8ab7b7341601d1141e51c3f5edb5351b 100644 --- a/src/KOKKOS/pair_lj_charmm_coul_long_kokkos.cpp +++ b/src/KOKKOS/pair_lj_charmm_coul_long_kokkos.cpp @@ -31,7 +31,7 @@ #include "integrate.h" #include "respa.h" #include "math_const.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -73,8 +73,8 @@ template<class DeviceType> PairLJCharmmCoulLongKokkos<DeviceType>::~PairLJCharmmCoulLongKokkos() { if (!copymode) { - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); k_cutsq = DAT::tdual_ffloat_2d(); k_cut_ljsq = DAT::tdual_ffloat_2d(); k_cut_coulsq = DAT::tdual_ffloat_2d(); @@ -118,13 +118,13 @@ void PairLJCharmmCoulLongKokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.view<DeviceType>(); } @@ -336,14 +336,14 @@ void PairLJCharmmCoulLongKokkos<DeviceType>::allocate() int n = atom->ntypes; memory->destroy(cutsq); - memory->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); + memoryKK->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); d_cutsq = k_cutsq.template view<DeviceType>(); //memory->destroy(cut_ljsq); - memory->create_kokkos(k_cut_ljsq,n+1,n+1,"pair:cut_ljsq"); + memoryKK->create_kokkos(k_cut_ljsq,n+1,n+1,"pair:cut_ljsq"); d_cut_ljsq = k_cut_ljsq.template view<DeviceType>(); - memory->create_kokkos(k_cut_coulsq,n+1,n+1,"pair:cut_coulsq"); + memoryKK->create_kokkos(k_cut_coulsq,n+1,n+1,"pair:cut_coulsq"); d_cut_coulsq = k_cut_coulsq.template view<DeviceType>(); k_params = Kokkos::DualView<params_lj_coul**,Kokkos::LayoutRight,DeviceType>("PairLJCharmmCoulLong::params",n+1,n+1); diff --git a/src/KOKKOS/pair_lj_class2_coul_cut_kokkos.cpp b/src/KOKKOS/pair_lj_class2_coul_cut_kokkos.cpp index 0081aca4f15ddf196ceb6d104ce75798fa55c736..3f355b1b160bf34df464b833cf1c938e76f5a17e 100644 --- a/src/KOKKOS/pair_lj_class2_coul_cut_kokkos.cpp +++ b/src/KOKKOS/pair_lj_class2_coul_cut_kokkos.cpp @@ -27,7 +27,7 @@ #include "integrate.h" #include "respa.h" #include "math_const.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -60,9 +60,9 @@ template<class DeviceType> PairLJClass2CoulCutKokkos<DeviceType>::~PairLJClass2CoulCutKokkos() { if (!copymode) { - memory->destroy_kokkos(k_cutsq, cutsq); - memory->destroy_kokkos(k_cut_ljsq, cut_ljsq); - memory->destroy_kokkos(k_cut_coulsq, cut_coulsq); + memoryKK->destroy_kokkos(k_cutsq, cutsq); + memoryKK->destroy_kokkos(k_cut_ljsq, cut_ljsq); + memoryKK->destroy_kokkos(k_cut_coulsq, cut_coulsq); } } @@ -95,13 +95,13 @@ void PairLJClass2CoulCutKokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.view<DeviceType>(); } @@ -255,13 +255,13 @@ void PairLJClass2CoulCutKokkos<DeviceType>::allocate() int n = atom->ntypes; memory->destroy(cutsq); - memory->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); + memoryKK->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); d_cutsq = k_cutsq.template view<DeviceType>(); memory->destroy(cut_ljsq); - memory->create_kokkos(k_cut_ljsq,cut_ljsq,n+1,n+1,"pair:cut_ljsq"); + memoryKK->create_kokkos(k_cut_ljsq,cut_ljsq,n+1,n+1,"pair:cut_ljsq"); d_cut_ljsq = k_cut_ljsq.template view<DeviceType>(); memory->destroy(cut_coulsq); - memory->create_kokkos(k_cut_coulsq,cut_coulsq,n+1,n+1,"pair:cut_coulsq"); + memoryKK->create_kokkos(k_cut_coulsq,cut_coulsq,n+1,n+1,"pair:cut_coulsq"); d_cut_coulsq = k_cut_coulsq.template view<DeviceType>(); k_params = Kokkos::DualView<params_lj_coul**,Kokkos::LayoutRight,DeviceType>("PairLJClass2CoulCut::params",n+1,n+1); params = k_params.template view<DeviceType>(); diff --git a/src/KOKKOS/pair_lj_class2_coul_long_kokkos.cpp b/src/KOKKOS/pair_lj_class2_coul_long_kokkos.cpp index b5dc358feb3fed379c7236e164a0f15d0104e60b..47976ec610aecdef5d6c41bd9d612533e6bc16aa 100644 --- a/src/KOKKOS/pair_lj_class2_coul_long_kokkos.cpp +++ b/src/KOKKOS/pair_lj_class2_coul_long_kokkos.cpp @@ -27,7 +27,7 @@ #include "integrate.h" #include "respa.h" #include "math_const.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -69,8 +69,8 @@ template<class DeviceType> PairLJClass2CoulLongKokkos<DeviceType>::~PairLJClass2CoulLongKokkos() { if (!copymode){ - memory->destroy_kokkos(k_cutsq, cutsq); - memory->destroy_kokkos(k_cut_ljsq, cut_ljsq); + memoryKK->destroy_kokkos(k_cutsq, cutsq); + memoryKK->destroy_kokkos(k_cut_ljsq, cut_ljsq); } } @@ -103,13 +103,13 @@ void PairLJClass2CoulLongKokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.view<DeviceType>(); } @@ -307,13 +307,13 @@ void PairLJClass2CoulLongKokkos<DeviceType>::allocate() int n = atom->ntypes; memory->destroy(cutsq); - memory->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); + memoryKK->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); d_cutsq = k_cutsq.template view<DeviceType>(); memory->destroy(cut_ljsq); - memory->create_kokkos(k_cut_ljsq,cut_ljsq,n+1,n+1,"pair:cut_ljsq"); + memoryKK->create_kokkos(k_cut_ljsq,cut_ljsq,n+1,n+1,"pair:cut_ljsq"); d_cut_ljsq = k_cut_ljsq.template view<DeviceType>(); - memory->create_kokkos(k_cut_coulsq,n+1,n+1,"pair:cut_coulsq"); + memoryKK->create_kokkos(k_cut_coulsq,n+1,n+1,"pair:cut_coulsq"); d_cut_coulsq = k_cut_coulsq.template view<DeviceType>(); k_params = Kokkos::DualView<params_lj_coul**,Kokkos::LayoutRight,DeviceType>("PairLJClass2CoulLong::params",n+1,n+1); params = k_params.template view<DeviceType>(); diff --git a/src/KOKKOS/pair_lj_class2_kokkos.cpp b/src/KOKKOS/pair_lj_class2_kokkos.cpp index 34cc15279b4d7a2708fc5c537d14ea44fe4e0dca..5beb520c00feb1971f5248e4ab31a7442133faa7 100644 --- a/src/KOKKOS/pair_lj_class2_kokkos.cpp +++ b/src/KOKKOS/pair_lj_class2_kokkos.cpp @@ -31,7 +31,7 @@ #include "integrate.h" #include "respa.h" #include "math_const.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -95,13 +95,13 @@ void PairLJClass2Kokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.view<DeviceType>(); } @@ -200,7 +200,7 @@ void PairLJClass2Kokkos<DeviceType>::allocate() int n = atom->ntypes; memory->destroy(cutsq); - memory->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); + memoryKK->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); d_cutsq = k_cutsq.template view<DeviceType>(); k_params = Kokkos::DualView<params_lj**,Kokkos::LayoutRight,DeviceType>("PairLJClass2::params",n+1,n+1); params = k_params.template view<DeviceType>(); diff --git a/src/KOKKOS/pair_lj_cut_coul_cut_kokkos.cpp b/src/KOKKOS/pair_lj_cut_coul_cut_kokkos.cpp index c3fda01ce63244bfa591deb80c23e4327288c68a..86e37bea77a6ffadf9ab082a36005ebca6da8aba 100644 --- a/src/KOKKOS/pair_lj_cut_coul_cut_kokkos.cpp +++ b/src/KOKKOS/pair_lj_cut_coul_cut_kokkos.cpp @@ -27,7 +27,7 @@ #include "integrate.h" #include "respa.h" #include "math_const.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -60,9 +60,9 @@ template<class DeviceType> PairLJCutCoulCutKokkos<DeviceType>::~PairLJCutCoulCutKokkos() { if (allocated){ - memory->destroy_kokkos(k_cutsq, cutsq); - memory->destroy_kokkos(k_cut_ljsq, cut_ljsq); - memory->destroy_kokkos(k_cut_coulsq, cut_coulsq); + memoryKK->destroy_kokkos(k_cutsq, cutsq); + memoryKK->destroy_kokkos(k_cut_ljsq, cut_ljsq); + memoryKK->destroy_kokkos(k_cut_coulsq, cut_coulsq); } } @@ -95,13 +95,13 @@ void PairLJCutCoulCutKokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.view<DeviceType>(); } @@ -247,13 +247,13 @@ void PairLJCutCoulCutKokkos<DeviceType>::allocate() int n = atom->ntypes; memory->destroy(cutsq); - memory->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); + memoryKK->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); d_cutsq = k_cutsq.template view<DeviceType>(); memory->destroy(cut_ljsq); - memory->create_kokkos(k_cut_ljsq,cut_ljsq,n+1,n+1,"pair:cut_ljsq"); + memoryKK->create_kokkos(k_cut_ljsq,cut_ljsq,n+1,n+1,"pair:cut_ljsq"); d_cut_ljsq = k_cut_ljsq.template view<DeviceType>(); memory->destroy(cut_coulsq); - memory->create_kokkos(k_cut_coulsq,cut_coulsq,n+1,n+1,"pair:cut_coulsq"); + memoryKK->create_kokkos(k_cut_coulsq,cut_coulsq,n+1,n+1,"pair:cut_coulsq"); d_cut_coulsq = k_cut_coulsq.template view<DeviceType>(); k_params = Kokkos::DualView<params_lj_coul**,Kokkos::LayoutRight,DeviceType>("PairLJCutCoulCut::params",n+1,n+1); params = k_params.template view<DeviceType>(); diff --git a/src/KOKKOS/pair_lj_cut_coul_debye_kokkos.cpp b/src/KOKKOS/pair_lj_cut_coul_debye_kokkos.cpp index b7a71cb99ac3d7f22ace36d5537a1dcb920dd60b..76e488c68667e4ac7b19150ba70e876ebc81fe7e 100644 --- a/src/KOKKOS/pair_lj_cut_coul_debye_kokkos.cpp +++ b/src/KOKKOS/pair_lj_cut_coul_debye_kokkos.cpp @@ -31,7 +31,7 @@ #include "integrate.h" #include "respa.h" #include "math_const.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -64,9 +64,9 @@ template<class DeviceType> PairLJCutCoulDebyeKokkos<DeviceType>::~PairLJCutCoulDebyeKokkos() { if (!copymode) { - memory->destroy_kokkos(k_cutsq, cutsq); - memory->destroy_kokkos(k_cut_ljsq, cut_ljsq); - memory->destroy_kokkos(k_cut_coulsq, cut_coulsq); + memoryKK->destroy_kokkos(k_cutsq, cutsq); + memoryKK->destroy_kokkos(k_cut_ljsq, cut_ljsq); + memoryKK->destroy_kokkos(k_cut_coulsq, cut_coulsq); } } @@ -99,13 +99,13 @@ void PairLJCutCoulDebyeKokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.view<DeviceType>(); } @@ -261,13 +261,13 @@ void PairLJCutCoulDebyeKokkos<DeviceType>::allocate() int n = atom->ntypes; memory->destroy(cutsq); - memory->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); + memoryKK->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); d_cutsq = k_cutsq.template view<DeviceType>(); memory->destroy(cut_ljsq); - memory->create_kokkos(k_cut_ljsq,cut_ljsq,n+1,n+1,"pair:cut_ljsq"); + memoryKK->create_kokkos(k_cut_ljsq,cut_ljsq,n+1,n+1,"pair:cut_ljsq"); d_cut_ljsq = k_cut_ljsq.template view<DeviceType>(); memory->destroy(cut_coulsq); - memory->create_kokkos(k_cut_coulsq,cut_coulsq,n+1,n+1,"pair:cut_coulsq"); + memoryKK->create_kokkos(k_cut_coulsq,cut_coulsq,n+1,n+1,"pair:cut_coulsq"); d_cut_coulsq = k_cut_coulsq.template view<DeviceType>(); k_params = Kokkos::DualView<params_lj_coul**,Kokkos::LayoutRight,DeviceType>("PairLJCutCoulDebye::params",n+1,n+1); params = k_params.template view<DeviceType>(); diff --git a/src/KOKKOS/pair_lj_cut_coul_dsf_kokkos.cpp b/src/KOKKOS/pair_lj_cut_coul_dsf_kokkos.cpp index 9df59636769a25dc3d512f36d4912486cb43be66..0da5e7f5f511432e5b8e81036caed9ec779dc992 100644 --- a/src/KOKKOS/pair_lj_cut_coul_dsf_kokkos.cpp +++ b/src/KOKKOS/pair_lj_cut_coul_dsf_kokkos.cpp @@ -31,7 +31,7 @@ #include "integrate.h" #include "respa.h" #include "math_const.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -72,9 +72,9 @@ template<class DeviceType> PairLJCutCoulDSFKokkos<DeviceType>::~PairLJCutCoulDSFKokkos() { if (!copymode) { - memory->destroy_kokkos(k_cutsq, cutsq); - memory->destroy_kokkos(k_cut_ljsq, cut_ljsq); - //memory->destroy_kokkos(k_cut_coulsq, cut_coulsq); + memoryKK->destroy_kokkos(k_cutsq, cutsq); + memoryKK->destroy_kokkos(k_cut_ljsq, cut_ljsq); + //memoryKK->destroy_kokkos(k_cut_coulsq, cut_coulsq); } } @@ -107,13 +107,13 @@ void PairLJCutCoulDSFKokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.view<DeviceType>(); } @@ -279,13 +279,13 @@ void PairLJCutCoulDSFKokkos<DeviceType>::allocate() int n = atom->ntypes; memory->destroy(cutsq); - memory->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); + memoryKK->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); d_cutsq = k_cutsq.template view<DeviceType>(); memory->destroy(cut_ljsq); - memory->create_kokkos(k_cut_ljsq,cut_ljsq,n+1,n+1,"pair:cut_ljsq"); + memoryKK->create_kokkos(k_cut_ljsq,cut_ljsq,n+1,n+1,"pair:cut_ljsq"); d_cut_ljsq = k_cut_ljsq.template view<DeviceType>(); //memory->destroy(cut_coulsq); - memory->create_kokkos(k_cut_coulsq,n+1,n+1,"pair:cut_coulsq"); + memoryKK->create_kokkos(k_cut_coulsq,n+1,n+1,"pair:cut_coulsq"); d_cut_coulsq = k_cut_coulsq.template view<DeviceType>(); k_params = Kokkos::DualView<params_lj_coul**,Kokkos::LayoutRight,DeviceType>("PairLJCutCoulDSF::params",n+1,n+1); params = k_params.template view<DeviceType>(); diff --git a/src/KOKKOS/pair_lj_cut_coul_long_kokkos.cpp b/src/KOKKOS/pair_lj_cut_coul_long_kokkos.cpp index 9bd79c7341a7116d74eb789ace8f5a984564278c..77a604534e2ae907a16db82ed013c26af2ec1042 100644 --- a/src/KOKKOS/pair_lj_cut_coul_long_kokkos.cpp +++ b/src/KOKKOS/pair_lj_cut_coul_long_kokkos.cpp @@ -27,7 +27,7 @@ #include "integrate.h" #include "respa.h" #include "math_const.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -68,13 +68,13 @@ PairLJCutCoulLongKokkos<DeviceType>::PairLJCutCoulLongKokkos(LAMMPS *lmp):PairLJ template<class DeviceType> PairLJCutCoulLongKokkos<DeviceType>::~PairLJCutCoulLongKokkos() { - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); eatom = NULL; vatom = NULL; if (allocated){ - memory->destroy_kokkos(k_cutsq, cutsq); - memory->destroy_kokkos(k_cut_ljsq, cut_ljsq); + memoryKK->destroy_kokkos(k_cutsq, cutsq); + memoryKK->destroy_kokkos(k_cut_ljsq, cut_ljsq); } } @@ -108,13 +108,13 @@ void PairLJCutCoulLongKokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.view<DeviceType>(); } @@ -304,13 +304,13 @@ void PairLJCutCoulLongKokkos<DeviceType>::allocate() int n = atom->ntypes; memory->destroy(cutsq); - memory->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); + memoryKK->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); d_cutsq = k_cutsq.template view<DeviceType>(); memory->destroy(cut_ljsq); - memory->create_kokkos(k_cut_ljsq,cut_ljsq,n+1,n+1,"pair:cut_ljsq"); + memoryKK->create_kokkos(k_cut_ljsq,cut_ljsq,n+1,n+1,"pair:cut_ljsq"); d_cut_ljsq = k_cut_ljsq.template view<DeviceType>(); - memory->create_kokkos(k_cut_coulsq,n+1,n+1,"pair:cut_coulsq"); + memoryKK->create_kokkos(k_cut_coulsq,n+1,n+1,"pair:cut_coulsq"); d_cut_coulsq = k_cut_coulsq.template view<DeviceType>(); k_params = Kokkos::DualView<params_lj_coul**,Kokkos::LayoutRight,DeviceType>("PairLJCutCoulLong::params",n+1,n+1); params = k_params.template view<DeviceType>(); diff --git a/src/KOKKOS/pair_lj_cut_kokkos.cpp b/src/KOKKOS/pair_lj_cut_kokkos.cpp index c90d5ad11c5295083eaaa4516c6f108665280def..2a228fb168be97184fda12d80b012745b286fa80 100644 --- a/src/KOKKOS/pair_lj_cut_kokkos.cpp +++ b/src/KOKKOS/pair_lj_cut_kokkos.cpp @@ -27,7 +27,7 @@ #include "integrate.h" #include "respa.h" #include "math_const.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -57,8 +57,8 @@ template<class DeviceType> PairLJCutKokkos<DeviceType>::~PairLJCutKokkos() { if (allocated) { - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); k_cutsq = DAT::tdual_ffloat_2d(); memory->sfree(cutsq); eatom = NULL; @@ -95,13 +95,13 @@ void PairLJCutKokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.view<DeviceType>(); } @@ -195,7 +195,7 @@ void PairLJCutKokkos<DeviceType>::allocate() int n = atom->ntypes; memory->destroy(cutsq); - memory->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); + memoryKK->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); d_cutsq = k_cutsq.template view<DeviceType>(); k_params = Kokkos::DualView<params_lj**,Kokkos::LayoutRight,DeviceType>("PairLJCut::params",n+1,n+1); params = k_params.template view<DeviceType>(); diff --git a/src/KOKKOS/pair_lj_expand_kokkos.cpp b/src/KOKKOS/pair_lj_expand_kokkos.cpp index 95ec252ad5fdebf3d4829abc524927d17ffa6eef..aca7202b78be66a04ce7443e7eb4f34c34a6a6e0 100644 --- a/src/KOKKOS/pair_lj_expand_kokkos.cpp +++ b/src/KOKKOS/pair_lj_expand_kokkos.cpp @@ -31,7 +31,7 @@ #include "integrate.h" #include "respa.h" #include "math_const.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -94,13 +94,13 @@ void PairLJExpandKokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.view<DeviceType>(); } @@ -202,7 +202,7 @@ void PairLJExpandKokkos<DeviceType>::allocate() int n = atom->ntypes; memory->destroy(cutsq); - memory->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); + memoryKK->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); d_cutsq = k_cutsq.template view<DeviceType>(); k_params = Kokkos::DualView<params_lj**,Kokkos::LayoutRight,DeviceType>("PairLJExpand::params",n+1,n+1); params = k_params.template view<DeviceType>(); diff --git a/src/KOKKOS/pair_lj_gromacs_coul_gromacs_kokkos.cpp b/src/KOKKOS/pair_lj_gromacs_coul_gromacs_kokkos.cpp index 4b21b08eb352ad95c8806bff4ae3175204eaf767..bf269288e076b3c69edf2cbff706426740a5e9f6 100644 --- a/src/KOKKOS/pair_lj_gromacs_coul_gromacs_kokkos.cpp +++ b/src/KOKKOS/pair_lj_gromacs_coul_gromacs_kokkos.cpp @@ -31,7 +31,7 @@ #include "integrate.h" #include "respa.h" #include "math_const.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -64,8 +64,8 @@ template<class DeviceType> PairLJGromacsCoulGromacsKokkos<DeviceType>::~PairLJGromacsCoulGromacsKokkos() { if (!copymode) { - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); k_cutsq = DAT::tdual_ffloat_2d(); k_cut_ljsq = DAT::tdual_ffloat_2d(); k_cut_coulsq = DAT::tdual_ffloat_2d(); @@ -109,13 +109,13 @@ void PairLJGromacsCoulGromacsKokkos<DeviceType>::compute(int eflag_in, int vflag // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.view<DeviceType>(); } @@ -299,14 +299,14 @@ void PairLJGromacsCoulGromacsKokkos<DeviceType>::allocate() int n = atom->ntypes; memory->destroy(cutsq); - memory->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); + memoryKK->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); d_cutsq = k_cutsq.template view<DeviceType>(); //memory->destroy(cut_ljsq); - memory->create_kokkos(k_cut_ljsq,n+1,n+1,"pair:cut_ljsq"); + memoryKK->create_kokkos(k_cut_ljsq,n+1,n+1,"pair:cut_ljsq"); d_cut_ljsq = k_cut_ljsq.template view<DeviceType>(); - memory->create_kokkos(k_cut_coulsq,n+1,n+1,"pair:cut_coulsq"); + memoryKK->create_kokkos(k_cut_coulsq,n+1,n+1,"pair:cut_coulsq"); d_cut_coulsq = k_cut_coulsq.template view<DeviceType>(); k_params = Kokkos::DualView<params_lj_coul_gromacs**,Kokkos::LayoutRight,DeviceType>("PairLJGromacsCoulGromacs::params",n+1,n+1); diff --git a/src/KOKKOS/pair_lj_gromacs_kokkos.cpp b/src/KOKKOS/pair_lj_gromacs_kokkos.cpp index 73a17d7b22f577fc31e4d6f3271ff4bab1b7694a..f24ff718c1042992bf47e18a6ce12e54ec7c169f 100644 --- a/src/KOKKOS/pair_lj_gromacs_kokkos.cpp +++ b/src/KOKKOS/pair_lj_gromacs_kokkos.cpp @@ -31,7 +31,7 @@ #include "integrate.h" #include "respa.h" #include "math_const.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -64,8 +64,8 @@ template<class DeviceType> PairLJGromacsKokkos<DeviceType>::~PairLJGromacsKokkos() { if (!copymode) { - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); k_cutsq = DAT::tdual_ffloat_2d(); k_cut_inner_sq = DAT::tdual_ffloat_2d(); memory->sfree(cutsq); @@ -106,13 +106,13 @@ void PairLJGromacsKokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.view<DeviceType>(); } @@ -240,15 +240,15 @@ void PairLJGromacsKokkos<DeviceType>::allocate() int n = atom->ntypes; memory->destroy(cutsq); - memory->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); + memoryKK->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); d_cutsq = k_cutsq.template view<DeviceType>(); memory->destroy(cut_inner); - memory->create_kokkos(k_cut_inner,cut_inner,n+1,n+1,"pair:cut_inner"); + memoryKK->create_kokkos(k_cut_inner,cut_inner,n+1,n+1,"pair:cut_inner"); d_cut_inner = k_cut_inner.template view<DeviceType>(); memory->destroy(cut_inner_sq); - memory->create_kokkos(k_cut_inner_sq,cut_inner_sq,n+1,n+1,"pair:cut_inner_sq"); + memoryKK->create_kokkos(k_cut_inner_sq,cut_inner_sq,n+1,n+1,"pair:cut_inner_sq"); d_cut_inner_sq = k_cut_inner_sq.template view<DeviceType>(); k_params = Kokkos::DualView<params_lj**,Kokkos::LayoutRight,DeviceType>("PairLJGromacs::params",n+1,n+1); diff --git a/src/KOKKOS/pair_lj_sdk_kokkos.cpp b/src/KOKKOS/pair_lj_sdk_kokkos.cpp index 2063f62b203546f4dfa18fd505b8e64ed7148968..aa579d5dc12995475731589da686ee4271867db9 100644 --- a/src/KOKKOS/pair_lj_sdk_kokkos.cpp +++ b/src/KOKKOS/pair_lj_sdk_kokkos.cpp @@ -27,7 +27,7 @@ #include "integrate.h" #include "respa.h" #include "math_const.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -94,13 +94,13 @@ void PairLJSDKKokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.view<DeviceType>(); } @@ -231,7 +231,7 @@ void PairLJSDKKokkos<DeviceType>::allocate() int n = atom->ntypes; memory->destroy(cutsq); - memory->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); + memoryKK->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); d_cutsq = k_cutsq.template view<DeviceType>(); k_params = Kokkos::DualView<params_lj**,Kokkos::LayoutRight,DeviceType>("PairLJSDK::params",n+1,n+1); params = k_params.template view<DeviceType>(); diff --git a/src/KOKKOS/pair_morse_kokkos.cpp b/src/KOKKOS/pair_morse_kokkos.cpp index 08a9b82640b4ea439c67457210d6baf31cb34027..5768d7e42cedba7cbdfb16bec74a765af657b837 100644 --- a/src/KOKKOS/pair_morse_kokkos.cpp +++ b/src/KOKKOS/pair_morse_kokkos.cpp @@ -31,7 +31,7 @@ #include "integrate.h" #include "respa.h" #include "math_const.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -61,8 +61,8 @@ template<class DeviceType> PairMorseKokkos<DeviceType>::~PairMorseKokkos() { if (allocated) { - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); k_cutsq = DAT::tdual_ffloat_2d(); memory->sfree(cutsq); eatom = NULL; @@ -99,13 +99,13 @@ void PairMorseKokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.view<DeviceType>(); } @@ -209,7 +209,7 @@ void PairMorseKokkos<DeviceType>::allocate() int n = atom->ntypes; memory->destroy(cutsq); - memory->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); + memoryKK->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq"); d_cutsq = k_cutsq.template view<DeviceType>(); k_params = Kokkos::DualView<params_morse**,Kokkos::LayoutRight,DeviceType>("PairMorse::params",n+1,n+1); params = k_params.template view<DeviceType>(); diff --git a/src/KOKKOS/pair_multi_lucy_rx_kokkos.cpp b/src/KOKKOS/pair_multi_lucy_rx_kokkos.cpp index d9a4f1ab83643ebd27f692c6290e436079c91522..0961cf44eb09ea1211ab7898c9c6532cbd518a14 100644 --- a/src/KOKKOS/pair_multi_lucy_rx_kokkos.cpp +++ b/src/KOKKOS/pair_multi_lucy_rx_kokkos.cpp @@ -31,13 +31,14 @@ #include "force.h" #include "comm.h" #include "neigh_list.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "citeme.h" #include "modify.h" #include "fix.h" #include "atom_masks.h" #include "neigh_request.h" +#include "kokkos.h" using namespace LAMMPS_NS; @@ -80,10 +81,10 @@ PairMultiLucyRXKokkos<DeviceType>::~PairMultiLucyRXKokkos() { if (copymode) return; - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); - memory->destroy_kokkos(k_cutsq,cutsq); + memoryKK->destroy_kokkos(k_cutsq,cutsq); delete h_table; delete d_table; @@ -153,13 +154,13 @@ void PairMultiLucyRXKokkos<DeviceType>::compute_style(int eflag_in, int vflag_in // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.template view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.template view<DeviceType>(); } @@ -864,20 +865,20 @@ void PairMultiLucyRXKokkos<DeviceType>::create_kokkos_tables() { const int tlm1 = tablength-1; - memory->create_kokkos(d_table->innersq,h_table->innersq,ntables,"Table::innersq"); - memory->create_kokkos(d_table->invdelta,h_table->invdelta,ntables,"Table::invdelta"); + memoryKK->create_kokkos(d_table->innersq,h_table->innersq,ntables,"Table::innersq"); + memoryKK->create_kokkos(d_table->invdelta,h_table->invdelta,ntables,"Table::invdelta"); if(tabstyle == LOOKUP) { - memory->create_kokkos(d_table->e,h_table->e,ntables,tlm1,"Table::e"); - memory->create_kokkos(d_table->f,h_table->f,ntables,tlm1,"Table::f"); + memoryKK->create_kokkos(d_table->e,h_table->e,ntables,tlm1,"Table::e"); + memoryKK->create_kokkos(d_table->f,h_table->f,ntables,tlm1,"Table::f"); } if(tabstyle == LINEAR) { - memory->create_kokkos(d_table->rsq,h_table->rsq,ntables,tablength,"Table::rsq"); - memory->create_kokkos(d_table->e,h_table->e,ntables,tablength,"Table::e"); - memory->create_kokkos(d_table->f,h_table->f,ntables,tablength,"Table::f"); - memory->create_kokkos(d_table->de,h_table->de,ntables,tlm1,"Table::de"); - memory->create_kokkos(d_table->df,h_table->df,ntables,tlm1,"Table::df"); + memoryKK->create_kokkos(d_table->rsq,h_table->rsq,ntables,tablength,"Table::rsq"); + memoryKK->create_kokkos(d_table->e,h_table->e,ntables,tablength,"Table::e"); + memoryKK->create_kokkos(d_table->f,h_table->f,ntables,tablength,"Table::f"); + memoryKK->create_kokkos(d_table->de,h_table->de,ntables,tlm1,"Table::de"); + memoryKK->create_kokkos(d_table->df,h_table->df,ntables,tlm1,"Table::df"); } for(int i=0; i < ntables; i++) { @@ -931,11 +932,11 @@ void PairMultiLucyRXKokkos<DeviceType>::allocate() memory->create(setflag,nt,nt,"pair:setflag"); - memory->create_kokkos(k_cutsq,cutsq,nt,nt,"pair:cutsq"); + memoryKK->create_kokkos(k_cutsq,cutsq,nt,nt,"pair:cutsq"); d_cutsq = k_cutsq.template view<DeviceType>(); k_cutsq.template modify<LMPHostType>(); - memory->create_kokkos(d_table->tabindex,h_table->tabindex,tabindex,nt,nt,"pair:tabindex"); + memoryKK->create_kokkos(d_table->tabindex,h_table->tabindex,tabindex,nt,nt,"pair:tabindex"); d_table_const.tabindex = d_table->tabindex; memset(&setflag[0][0],0,nt*nt*sizeof(int)); diff --git a/src/KOKKOS/pair_multi_lucy_rx_kokkos.h b/src/KOKKOS/pair_multi_lucy_rx_kokkos.h index b8ced4c847f6088efe98c09f50ea706626531883..aee1763b06e25c019a996c8b1c0c2b5df399188a 100644 --- a/src/KOKKOS/pair_multi_lucy_rx_kokkos.h +++ b/src/KOKKOS/pair_multi_lucy_rx_kokkos.h @@ -25,6 +25,7 @@ PairStyle(multi/lucy/rx/kk/host,PairMultiLucyRXKokkos<LMPHostType>) #include "pair_multi_lucy_rx.h" #include "pair_kokkos.h" +#include "kokkos_base.h" #include "kokkos_type.h" namespace LAMMPS_NS { @@ -43,7 +44,7 @@ template<int NEIGHFLAG, int NEWTON_PAIR, bool ONE_TYPE> struct TagPairMultiLucyRXComputeLocalDensity{}; template<class DeviceType> -class PairMultiLucyRXKokkos : public PairMultiLucyRX { +class PairMultiLucyRXKokkos : public PairMultiLucyRX, public KokkosBase { public: typedef DeviceType device_type; typedef ArrayTypes<DeviceType> AT; diff --git a/src/KOKKOS/pair_reaxc_kokkos.cpp b/src/KOKKOS/pair_reaxc_kokkos.cpp index 1131cb787fc2e5e1978803a39336a10900ac001f..1f596377679fcac252660addc37109675da0260e 100644 --- a/src/KOKKOS/pair_reaxc_kokkos.cpp +++ b/src/KOKKOS/pair_reaxc_kokkos.cpp @@ -32,12 +32,13 @@ #include "respa.h" #include "math_const.h" #include "math_special.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" #include "reaxc_defs.h" #include "reaxc_lookup.h" #include "reaxc_tool_box.h" +#include "modify.h" #define TEAMSIZE 128 @@ -81,12 +82,12 @@ PairReaxCKokkos<DeviceType>::~PairReaxCKokkos() { if (copymode) return; - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); - memory->destroy_kokkos(k_tmpid,tmpid); + memoryKK->destroy_kokkos(k_tmpid,tmpid); tmpid = NULL; - memory->destroy_kokkos(k_tmpbo,tmpbo); + memoryKK->destroy_kokkos(k_tmpbo,tmpbo); tmpbo = NULL; } @@ -1339,10 +1340,10 @@ void PairReaxCKokkos<DeviceType>::allocate_array() // FixReaxCSpecies if (fixspecies_flag) { - memory->destroy_kokkos(k_tmpid,tmpid); - memory->destroy_kokkos(k_tmpbo,tmpbo); - memory->create_kokkos(k_tmpid,tmpid,nmax,MAXSPECBOND,"pair:tmpid"); - memory->create_kokkos(k_tmpbo,tmpbo,nmax,MAXSPECBOND,"pair:tmpbo"); + memoryKK->destroy_kokkos(k_tmpid,tmpid); + memoryKK->destroy_kokkos(k_tmpbo,tmpbo); + memoryKK->create_kokkos(k_tmpid,tmpid,nmax,MAXSPECBOND,"pair:tmpid"); + memoryKK->create_kokkos(k_tmpbo,tmpbo,nmax,MAXSPECBOND,"pair:tmpbo"); } // FixReaxCBonds @@ -3911,14 +3912,14 @@ void PairReaxCKokkos<DeviceType>::ev_setup(int eflag, int vflag) if (eflag_atom && atom->nmax > maxeatom) { maxeatom = atom->nmax; - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); v_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom && atom->nmax > maxvatom) { maxvatom = atom->nmax; - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); v_vatom = k_vatom.view<DeviceType>(); } diff --git a/src/KOKKOS/pair_sw_kokkos.cpp b/src/KOKKOS/pair_sw_kokkos.cpp index 3440f7c63923da0467c61b702c1184ab3cb1856c..63b4c19ae1d7c40d4e3234562b6b7eb8c54e9470 100644 --- a/src/KOKKOS/pair_sw_kokkos.cpp +++ b/src/KOKKOS/pair_sw_kokkos.cpp @@ -27,10 +27,10 @@ #include "neigh_request.h" #include "force.h" #include "comm.h" -#include "memory.h" +#include "memory_kokkos.h" #include "neighbor.h" #include "neigh_list_kokkos.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" #include "math_const.h" @@ -63,8 +63,8 @@ template<class DeviceType> PairSWKokkos<DeviceType>::~PairSWKokkos() { if (!copymode) { - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); eatom = NULL; vatom = NULL; } @@ -86,13 +86,13 @@ void PairSWKokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.view<DeviceType>(); } diff --git a/src/KOKKOS/pair_table_kokkos.cpp b/src/KOKKOS/pair_table_kokkos.cpp index 7f763baae6fe7b40eb4830d243c0b714a6bcf2b1..b3e80d1f664e012555a826834df2d6b78ec7bba4 100644 --- a/src/KOKKOS/pair_table_kokkos.cpp +++ b/src/KOKKOS/pair_table_kokkos.cpp @@ -27,7 +27,7 @@ #include "neighbor.h" #include "neigh_list.h" #include "neigh_request.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -92,13 +92,13 @@ void PairTableKokkos<DeviceType>::compute_style(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.view<DeviceType>(); } @@ -265,41 +265,41 @@ void PairTableKokkos<DeviceType>::create_kokkos_tables() { const int tlm1 = tablength-1; - memory->create_kokkos(d_table->nshiftbits,h_table->nshiftbits,ntables,"Table::nshiftbits"); - memory->create_kokkos(d_table->nmask,h_table->nmask,ntables,"Table::nmask"); - memory->create_kokkos(d_table->innersq,h_table->innersq,ntables,"Table::innersq"); - memory->create_kokkos(d_table->invdelta,h_table->invdelta,ntables,"Table::invdelta"); - memory->create_kokkos(d_table->deltasq6,h_table->deltasq6,ntables,"Table::deltasq6"); + memoryKK->create_kokkos(d_table->nshiftbits,h_table->nshiftbits,ntables,"Table::nshiftbits"); + memoryKK->create_kokkos(d_table->nmask,h_table->nmask,ntables,"Table::nmask"); + memoryKK->create_kokkos(d_table->innersq,h_table->innersq,ntables,"Table::innersq"); + memoryKK->create_kokkos(d_table->invdelta,h_table->invdelta,ntables,"Table::invdelta"); + memoryKK->create_kokkos(d_table->deltasq6,h_table->deltasq6,ntables,"Table::deltasq6"); if(tabstyle == LOOKUP) { - memory->create_kokkos(d_table->e,h_table->e,ntables,tlm1,"Table::e"); - memory->create_kokkos(d_table->f,h_table->f,ntables,tlm1,"Table::f"); + memoryKK->create_kokkos(d_table->e,h_table->e,ntables,tlm1,"Table::e"); + memoryKK->create_kokkos(d_table->f,h_table->f,ntables,tlm1,"Table::f"); } if(tabstyle == LINEAR) { - memory->create_kokkos(d_table->rsq,h_table->rsq,ntables,tablength,"Table::rsq"); - memory->create_kokkos(d_table->e,h_table->e,ntables,tablength,"Table::e"); - memory->create_kokkos(d_table->f,h_table->f,ntables,tablength,"Table::f"); - memory->create_kokkos(d_table->de,h_table->de,ntables,tlm1,"Table::de"); - memory->create_kokkos(d_table->df,h_table->df,ntables,tlm1,"Table::df"); + memoryKK->create_kokkos(d_table->rsq,h_table->rsq,ntables,tablength,"Table::rsq"); + memoryKK->create_kokkos(d_table->e,h_table->e,ntables,tablength,"Table::e"); + memoryKK->create_kokkos(d_table->f,h_table->f,ntables,tablength,"Table::f"); + memoryKK->create_kokkos(d_table->de,h_table->de,ntables,tlm1,"Table::de"); + memoryKK->create_kokkos(d_table->df,h_table->df,ntables,tlm1,"Table::df"); } if(tabstyle == SPLINE) { - memory->create_kokkos(d_table->rsq,h_table->rsq,ntables,tablength,"Table::rsq"); - memory->create_kokkos(d_table->e,h_table->e,ntables,tablength,"Table::e"); - memory->create_kokkos(d_table->f,h_table->f,ntables,tablength,"Table::f"); - memory->create_kokkos(d_table->e2,h_table->e2,ntables,tablength,"Table::e2"); - memory->create_kokkos(d_table->f2,h_table->f2,ntables,tablength,"Table::f2"); + memoryKK->create_kokkos(d_table->rsq,h_table->rsq,ntables,tablength,"Table::rsq"); + memoryKK->create_kokkos(d_table->e,h_table->e,ntables,tablength,"Table::e"); + memoryKK->create_kokkos(d_table->f,h_table->f,ntables,tablength,"Table::f"); + memoryKK->create_kokkos(d_table->e2,h_table->e2,ntables,tablength,"Table::e2"); + memoryKK->create_kokkos(d_table->f2,h_table->f2,ntables,tablength,"Table::f2"); } if(tabstyle == BITMAP) { int ntable = 1 << tablength; - memory->create_kokkos(d_table->rsq,h_table->rsq,ntables,ntable,"Table::rsq"); - memory->create_kokkos(d_table->e,h_table->e,ntables,ntable,"Table::e"); - memory->create_kokkos(d_table->f,h_table->f,ntables,ntable,"Table::f"); - memory->create_kokkos(d_table->de,h_table->de,ntables,ntable,"Table::de"); - memory->create_kokkos(d_table->df,h_table->df,ntables,ntable,"Table::df"); - memory->create_kokkos(d_table->drsq,h_table->drsq,ntables,ntable,"Table::drsq"); + memoryKK->create_kokkos(d_table->rsq,h_table->rsq,ntables,ntable,"Table::rsq"); + memoryKK->create_kokkos(d_table->e,h_table->e,ntables,ntable,"Table::e"); + memoryKK->create_kokkos(d_table->f,h_table->f,ntables,ntable,"Table::f"); + memoryKK->create_kokkos(d_table->de,h_table->de,ntables,ntable,"Table::de"); + memoryKK->create_kokkos(d_table->df,h_table->df,ntables,ntable,"Table::df"); + memoryKK->create_kokkos(d_table->drsq,h_table->drsq,ntables,ntable,"Table::drsq"); } @@ -410,8 +410,8 @@ void PairTableKokkos<DeviceType>::allocate() const int nt = atom->ntypes + 1; memory->create(setflag,nt,nt,"pair:setflag"); - memory->create_kokkos(d_table->cutsq,h_table->cutsq,cutsq,nt,nt,"pair:cutsq"); - memory->create_kokkos(d_table->tabindex,h_table->tabindex,tabindex,nt,nt,"pair:tabindex"); + memoryKK->create_kokkos(d_table->cutsq,h_table->cutsq,cutsq,nt,nt,"pair:cutsq"); + memoryKK->create_kokkos(d_table->tabindex,h_table->tabindex,tabindex,nt,nt,"pair:tabindex"); d_table_const.cutsq = d_table->cutsq; d_table_const.tabindex = d_table->tabindex; diff --git a/src/KOKKOS/pair_table_rx_kokkos.cpp b/src/KOKKOS/pair_table_rx_kokkos.cpp index 2f5a670537ce05f528d1c789c3070292cacd8c80..7bc5198d8cef6e8bbfcf40ba24e4d0bb575b9cec 100644 --- a/src/KOKKOS/pair_table_rx_kokkos.cpp +++ b/src/KOKKOS/pair_table_rx_kokkos.cpp @@ -27,11 +27,13 @@ #include "neighbor.h" #include "neigh_list.h" #include "neigh_request.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" #include "fix.h" #include "kokkos_few.h" +#include "kokkos.h" +#include "modify.h" #include <cassert> using namespace LAMMPS_NS; @@ -162,12 +164,12 @@ PairTableRXKokkos<DeviceType>::~PairTableRXKokkos() delete [] site1; delete [] site2; - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); if (allocated) { - memory->destroy_kokkos(d_table->cutsq, cutsq); - memory->destroy_kokkos(d_table->tabindex, tabindex); + memoryKK->destroy_kokkos(d_table->cutsq, cutsq); + memoryKK->destroy_kokkos(d_table->tabindex, tabindex); } delete h_table; @@ -621,13 +623,13 @@ void PairTableRXKokkos<DeviceType>::compute_style(int eflag_in, int vflag_in) else evflag = vflag_fdotr = 0; if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.template view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.template view<DeviceType>(); } @@ -798,41 +800,41 @@ void PairTableRXKokkos<DeviceType>::create_kokkos_tables() { const int tlm1 = tablength-1; - memory->create_kokkos(d_table->nshiftbits,h_table->nshiftbits,ntables,"Table::nshiftbits"); - memory->create_kokkos(d_table->nmask,h_table->nmask,ntables,"Table::nmask"); - memory->create_kokkos(d_table->innersq,h_table->innersq,ntables,"Table::innersq"); - memory->create_kokkos(d_table->invdelta,h_table->invdelta,ntables,"Table::invdelta"); - memory->create_kokkos(d_table->deltasq6,h_table->deltasq6,ntables,"Table::deltasq6"); + memoryKK->create_kokkos(d_table->nshiftbits,h_table->nshiftbits,ntables,"Table::nshiftbits"); + memoryKK->create_kokkos(d_table->nmask,h_table->nmask,ntables,"Table::nmask"); + memoryKK->create_kokkos(d_table->innersq,h_table->innersq,ntables,"Table::innersq"); + memoryKK->create_kokkos(d_table->invdelta,h_table->invdelta,ntables,"Table::invdelta"); + memoryKK->create_kokkos(d_table->deltasq6,h_table->deltasq6,ntables,"Table::deltasq6"); if(tabstyle == LOOKUP) { - memory->create_kokkos(d_table->e,h_table->e,ntables,tlm1,"Table::e"); - memory->create_kokkos(d_table->f,h_table->f,ntables,tlm1,"Table::f"); + memoryKK->create_kokkos(d_table->e,h_table->e,ntables,tlm1,"Table::e"); + memoryKK->create_kokkos(d_table->f,h_table->f,ntables,tlm1,"Table::f"); } if(tabstyle == LINEAR) { - memory->create_kokkos(d_table->rsq,h_table->rsq,ntables,tablength,"Table::rsq"); - memory->create_kokkos(d_table->e,h_table->e,ntables,tablength,"Table::e"); - memory->create_kokkos(d_table->f,h_table->f,ntables,tablength,"Table::f"); - memory->create_kokkos(d_table->de,h_table->de,ntables,tlm1,"Table::de"); - memory->create_kokkos(d_table->df,h_table->df,ntables,tlm1,"Table::df"); + memoryKK->create_kokkos(d_table->rsq,h_table->rsq,ntables,tablength,"Table::rsq"); + memoryKK->create_kokkos(d_table->e,h_table->e,ntables,tablength,"Table::e"); + memoryKK->create_kokkos(d_table->f,h_table->f,ntables,tablength,"Table::f"); + memoryKK->create_kokkos(d_table->de,h_table->de,ntables,tlm1,"Table::de"); + memoryKK->create_kokkos(d_table->df,h_table->df,ntables,tlm1,"Table::df"); } if(tabstyle == SPLINE) { - memory->create_kokkos(d_table->rsq,h_table->rsq,ntables,tablength,"Table::rsq"); - memory->create_kokkos(d_table->e,h_table->e,ntables,tablength,"Table::e"); - memory->create_kokkos(d_table->f,h_table->f,ntables,tablength,"Table::f"); - memory->create_kokkos(d_table->e2,h_table->e2,ntables,tablength,"Table::e2"); - memory->create_kokkos(d_table->f2,h_table->f2,ntables,tablength,"Table::f2"); + memoryKK->create_kokkos(d_table->rsq,h_table->rsq,ntables,tablength,"Table::rsq"); + memoryKK->create_kokkos(d_table->e,h_table->e,ntables,tablength,"Table::e"); + memoryKK->create_kokkos(d_table->f,h_table->f,ntables,tablength,"Table::f"); + memoryKK->create_kokkos(d_table->e2,h_table->e2,ntables,tablength,"Table::e2"); + memoryKK->create_kokkos(d_table->f2,h_table->f2,ntables,tablength,"Table::f2"); } if(tabstyle == BITMAP) { int ntable = 1 << tablength; - memory->create_kokkos(d_table->rsq,h_table->rsq,ntables,ntable,"Table::rsq"); - memory->create_kokkos(d_table->e,h_table->e,ntables,ntable,"Table::e"); - memory->create_kokkos(d_table->f,h_table->f,ntables,ntable,"Table::f"); - memory->create_kokkos(d_table->de,h_table->de,ntables,ntable,"Table::de"); - memory->create_kokkos(d_table->df,h_table->df,ntables,ntable,"Table::df"); - memory->create_kokkos(d_table->drsq,h_table->drsq,ntables,ntable,"Table::drsq"); + memoryKK->create_kokkos(d_table->rsq,h_table->rsq,ntables,ntable,"Table::rsq"); + memoryKK->create_kokkos(d_table->e,h_table->e,ntables,ntable,"Table::e"); + memoryKK->create_kokkos(d_table->f,h_table->f,ntables,ntable,"Table::f"); + memoryKK->create_kokkos(d_table->de,h_table->de,ntables,ntable,"Table::de"); + memoryKK->create_kokkos(d_table->df,h_table->df,ntables,ntable,"Table::df"); + memoryKK->create_kokkos(d_table->drsq,h_table->drsq,ntables,ntable,"Table::drsq"); } @@ -943,8 +945,8 @@ void PairTableRXKokkos<DeviceType>::allocate() const int nt = atom->ntypes + 1; memory->create(setflag,nt,nt,"pair:setflag"); - memory->create_kokkos(d_table->cutsq,h_table->cutsq,cutsq,nt,nt,"pair:cutsq"); - memory->create_kokkos(d_table->tabindex,h_table->tabindex,tabindex,nt,nt,"pair:tabindex"); + memoryKK->create_kokkos(d_table->cutsq,h_table->cutsq,cutsq,nt,nt,"pair:cutsq"); + memoryKK->create_kokkos(d_table->tabindex,h_table->tabindex,tabindex,nt,nt,"pair:tabindex"); d_table_const.cutsq = d_table->cutsq; d_table_const.tabindex = d_table->tabindex; diff --git a/src/KOKKOS/pair_tersoff_kokkos.cpp b/src/KOKKOS/pair_tersoff_kokkos.cpp index 3a5c2227efe1b9f71e9b1277096882e86027336e..c585da60297b4a852b50ab19d00a060edaaf6f83 100644 --- a/src/KOKKOS/pair_tersoff_kokkos.cpp +++ b/src/KOKKOS/pair_tersoff_kokkos.cpp @@ -31,7 +31,7 @@ #include "integrate.h" #include "respa.h" #include "math_const.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -60,8 +60,8 @@ template<class DeviceType> PairTersoffKokkos<DeviceType>::~PairTersoffKokkos() { if (!copymode) { - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); } } @@ -170,13 +170,13 @@ void PairTersoffKokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.view<DeviceType>(); } diff --git a/src/KOKKOS/pair_tersoff_mod_kokkos.cpp b/src/KOKKOS/pair_tersoff_mod_kokkos.cpp index 9a59979f4c9d4d0279f5b830d0e66bcb1af39148..8e718fbf6f16df5017468b116c1b140a82eb0b4c 100644 --- a/src/KOKKOS/pair_tersoff_mod_kokkos.cpp +++ b/src/KOKKOS/pair_tersoff_mod_kokkos.cpp @@ -31,7 +31,7 @@ #include "integrate.h" #include "respa.h" #include "math_const.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -60,8 +60,8 @@ template<class DeviceType> PairTersoffMODKokkos<DeviceType>::~PairTersoffMODKokkos() { if (!copymode) { - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); } } @@ -170,13 +170,13 @@ void PairTersoffMODKokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.view<DeviceType>(); } diff --git a/src/KOKKOS/pair_tersoff_zbl_kokkos.cpp b/src/KOKKOS/pair_tersoff_zbl_kokkos.cpp index 8468bb01f5fc545d177ad637de70ae8514087693..11a5ff100753ec02dafa657dd481222bfaf7875c 100644 --- a/src/KOKKOS/pair_tersoff_zbl_kokkos.cpp +++ b/src/KOKKOS/pair_tersoff_zbl_kokkos.cpp @@ -31,7 +31,7 @@ #include "integrate.h" #include "respa.h" #include "math_const.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" @@ -71,8 +71,8 @@ template<class DeviceType> PairTersoffZBLKokkos<DeviceType>::~PairTersoffZBLKokkos() { if (!copymode) { - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); } } @@ -184,13 +184,13 @@ void PairTersoffZBLKokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.view<DeviceType>(); } diff --git a/src/KOKKOS/pair_vashishta_kokkos.cpp b/src/KOKKOS/pair_vashishta_kokkos.cpp index fe2394ae84a9fea6847c5f22eef598c25719b95a..e7dd01bfe4b48e903149db04ab8bb3b7af6cd3b6 100644 --- a/src/KOKKOS/pair_vashishta_kokkos.cpp +++ b/src/KOKKOS/pair_vashishta_kokkos.cpp @@ -27,10 +27,10 @@ #include "neigh_request.h" #include "force.h" #include "comm.h" -#include "memory.h" +#include "memory_kokkos.h" #include "neighbor.h" #include "neigh_list_kokkos.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" #include "math_const.h" @@ -62,8 +62,8 @@ template<class DeviceType> PairVashishtaKokkos<DeviceType>::~PairVashishtaKokkos() { if (!copymode) { - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); eatom = NULL; vatom = NULL; } @@ -85,13 +85,13 @@ void PairVashishtaKokkos<DeviceType>::compute(int eflag_in, int vflag_in) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.view<DeviceType>(); } diff --git a/src/KOKKOS/pppm_kokkos.cpp b/src/KOKKOS/pppm_kokkos.cpp index bd3ed3644f53373305a995382144a6d5c27a5432..cf6e2814c0e0bb56f5a19b2a2481ea019259e114 100644 --- a/src/KOKKOS/pppm_kokkos.cpp +++ b/src/KOKKOS/pppm_kokkos.cpp @@ -32,9 +32,10 @@ #include "domain.h" #include "fft3d_wrap.h" #include "remap_wrap.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" #include "atom_masks.h" +#include "kokkos.h" #include "math_const.h" #include "math_special_kokkos.h" @@ -162,8 +163,8 @@ PPPMKokkos<DeviceType>::~PPPMKokkos() //memory->destroy(part2grid); //memory->destroy(acons); - memory->destroy_kokkos(k_eatom,eatom); - memory->destroy_kokkos(k_vatom,vatom); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->destroy_kokkos(k_vatom,vatom); eatom = NULL; vatom = NULL; } @@ -618,13 +619,13 @@ void PPPMKokkos<DeviceType>::compute(int eflag, int vflag) // reallocate per-atom arrays if necessary if (eflag_atom) { - memory->destroy_kokkos(k_eatom,eatom); - memory->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); + memoryKK->destroy_kokkos(k_eatom,eatom); + memoryKK->create_kokkos(k_eatom,eatom,maxeatom,"pair:eatom"); d_eatom = k_eatom.view<DeviceType>(); } if (vflag_atom) { - memory->destroy_kokkos(k_vatom,vatom); - memory->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); + memoryKK->destroy_kokkos(k_vatom,vatom); + memoryKK->create_kokkos(k_vatom,vatom,maxvatom,6,"pair:vatom"); d_vatom = k_vatom.view<DeviceType>(); } @@ -805,12 +806,12 @@ void PPPMKokkos<DeviceType>::allocate() { d_density_brick = typename AT::t_FFT_SCALAR_3d("pppm:density_brick",nzhi_out-nzlo_out+1,nyhi_out-nylo_out+1,nxhi_out-nxlo_out+1); - memory->create_kokkos(k_density_fft,density_fft,nfft_both,"pppm:d_density_fft"); + memoryKK->create_kokkos(k_density_fft,density_fft,nfft_both,"pppm:d_density_fft"); d_density_fft = k_density_fft.view<DeviceType>(); d_greensfn = typename AT::t_float_1d("pppm:greensfn",nfft_both); - memory->create_kokkos(k_work1,work1,2*nfft_both,"pppm:work1"); - memory->create_kokkos(k_work2,work2,2*nfft_both,"pppm:work2"); + memoryKK->create_kokkos(k_work1,work1,2*nfft_both,"pppm:work1"); + memoryKK->create_kokkos(k_work2,work2,2*nfft_both,"pppm:work2"); d_work1 = k_work1.view<DeviceType>(); d_work2 = k_work2.view<DeviceType>(); d_vg = typename AT::t_virial_array("pppm:vg",nfft_both); @@ -878,13 +879,13 @@ void PPPMKokkos<DeviceType>::allocate() template<class DeviceType> void PPPMKokkos<DeviceType>::deallocate() { - memory->destroy_kokkos(d_density_fft,density_fft); + memoryKK->destroy_kokkos(d_density_fft,density_fft); density_fft = NULL; - memory->destroy_kokkos(d_greensfn,greensfn); + memoryKK->destroy_kokkos(d_greensfn,greensfn); greensfn = NULL; - memory->destroy_kokkos(d_work1,work1); + memoryKK->destroy_kokkos(d_work1,work1); work1 = NULL; - memory->destroy_kokkos(d_work2,work2); + memoryKK->destroy_kokkos(d_work2,work2); work2 = NULL; delete fft1; @@ -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 adec5ff1bd9d7520c2c469638d497759375c362a..5fa03a098925430d50f5dd09bc59612398a9d759 100644 --- a/src/KOKKOS/verlet_kokkos.cpp +++ b/src/KOKKOS/verlet_kokkos.cpp @@ -32,8 +32,9 @@ #include "compute.h" #include "fix.h" #include "timer.h" -#include "memory.h" +#include "memory_kokkos.h" #include "error.h" +#include "kokkos.h" #include <ctime> diff --git a/src/Purge.list b/src/Purge.list index 312994fdb7cf413464728e45dfec655a7b6feb95..93808b784ca58254f1c2076f47da5b773de3ea51 100644 --- a/src/Purge.list +++ b/src/Purge.list @@ -16,6 +16,9 @@ style_region.h style_neigh_bin.h style_neigh_pair.h style_neigh_stencil.h +# deleted on 1 December 2017 +npair_half_bin_newtoff_intel.cpp +npair_half_bin_newtoff_intel.h # deleted on 11 October 2017 fix_shear_history_omp.cpp fix_shear_history_omp.h diff --git a/src/USER-DRUDE/pair_lj_cut_thole_long.cpp b/src/USER-DRUDE/pair_lj_cut_thole_long.cpp index 4163a816ac9bcd8662fd25facb1200019071cd76..f78ced3d8bd96caaa35798db21f1bce5f19a05e2 100644 --- a/src/USER-DRUDE/pair_lj_cut_thole_long.cpp +++ b/src/USER-DRUDE/pair_lj_cut_thole_long.cpp @@ -32,6 +32,8 @@ #include "math_const.h" #include "memory.h" #include "error.h" +#include "modify.h" +#include "domain.h" using namespace LAMMPS_NS; using namespace MathConst; diff --git a/src/USER-DRUDE/pair_thole.cpp b/src/USER-DRUDE/pair_thole.cpp index abb37b82b766581aafadb5519f76169ca8289491..ace119ce523892b0ad7a8d52c0288f0bc0f725b4 100644 --- a/src/USER-DRUDE/pair_thole.cpp +++ b/src/USER-DRUDE/pair_thole.cpp @@ -25,6 +25,8 @@ #include "error.h" #include "fix.h" #include "fix_store.h" +#include "domain.h" +#include "modify.h" using namespace LAMMPS_NS; diff --git a/src/USER-INTEL/angle_charmm_intel.cpp b/src/USER-INTEL/angle_charmm_intel.cpp index 031c9642000c22a6aa7f4156db319de5409fce6e..bcaecb4696fa1d93e277d473e931b268298c5153 100644 --- a/src/USER-INTEL/angle_charmm_intel.cpp +++ b/src/USER-INTEL/angle_charmm_intel.cpp @@ -23,6 +23,7 @@ #include "domain.h" #include "comm.h" #include "force.h" +#include "modify.h" #include "math_const.h" #include "memory.h" #include "suffix.h" diff --git a/src/USER-INTEL/angle_harmonic_intel.cpp b/src/USER-INTEL/angle_harmonic_intel.cpp index 84220277d7a736131c30fa37d79fd93364e0c4ca..ffc81c496d77c654f35d6f0e1ad71a39a06844ee 100644 --- a/src/USER-INTEL/angle_harmonic_intel.cpp +++ b/src/USER-INTEL/angle_harmonic_intel.cpp @@ -23,6 +23,7 @@ #include "domain.h" #include "comm.h" #include "force.h" +#include "modify.h" #include "math_const.h" #include "memory.h" #include "suffix.h" diff --git a/src/USER-INTEL/bond_fene_intel.cpp b/src/USER-INTEL/bond_fene_intel.cpp index 93d64ed631f5466baec735c52a2770160b2eb8fb..004a2e5413d6d566cb536b9d4290e80540315e81 100644 --- a/src/USER-INTEL/bond_fene_intel.cpp +++ b/src/USER-INTEL/bond_fene_intel.cpp @@ -19,6 +19,7 @@ #include <stdlib.h> #include "bond_fene_intel.h" #include "atom.h" +#include "modify.h" #include "neighbor.h" #include "domain.h" #include "comm.h" diff --git a/src/USER-INTEL/bond_harmonic_intel.cpp b/src/USER-INTEL/bond_harmonic_intel.cpp index 0ac466f11386c99bef29b95a33e086349389239c..0c714edc0959519621b624d19dd6e79b5b081977 100644 --- a/src/USER-INTEL/bond_harmonic_intel.cpp +++ b/src/USER-INTEL/bond_harmonic_intel.cpp @@ -19,6 +19,7 @@ #include <stdlib.h> #include "bond_harmonic_intel.h" #include "atom.h" +#include "modify.h" #include "neighbor.h" #include "domain.h" #include "comm.h" diff --git a/src/USER-INTEL/dihedral_charmm_intel.cpp b/src/USER-INTEL/dihedral_charmm_intel.cpp index 0e13e92251594d04348d98f2200c3501a5b29fd5..c8429dc8ff7079930f00459c33ad85674ebaf222 100644 --- a/src/USER-INTEL/dihedral_charmm_intel.cpp +++ b/src/USER-INTEL/dihedral_charmm_intel.cpp @@ -21,6 +21,7 @@ #include "atom.h" #include "comm.h" #include "memory.h" +#include "modify.h" #include "neighbor.h" #include "domain.h" #include "force.h" diff --git a/src/USER-INTEL/dihedral_fourier_intel.cpp b/src/USER-INTEL/dihedral_fourier_intel.cpp index 805ffc0e256177f294bc8f2b118e778c4cd7054e..772ea5b02fdf4a84838bd7a59a90def0b8a3d107 100644 --- a/src/USER-INTEL/dihedral_fourier_intel.cpp +++ b/src/USER-INTEL/dihedral_fourier_intel.cpp @@ -21,6 +21,7 @@ #include "atom.h" #include "comm.h" #include "memory.h" +#include "modify.h" #include "neighbor.h" #include "domain.h" #include "force.h" diff --git a/src/USER-INTEL/dihedral_harmonic_intel.cpp b/src/USER-INTEL/dihedral_harmonic_intel.cpp index 5d16b0da745425abd66617ab50043ef796d78db3..b35ea4e03a9b67a3b619028c0c559b304cbbaf07 100644 --- a/src/USER-INTEL/dihedral_harmonic_intel.cpp +++ b/src/USER-INTEL/dihedral_harmonic_intel.cpp @@ -21,6 +21,7 @@ #include "atom.h" #include "comm.h" #include "memory.h" +#include "modify.h" #include "neighbor.h" #include "domain.h" #include "force.h" diff --git a/src/USER-INTEL/dihedral_opls_intel.cpp b/src/USER-INTEL/dihedral_opls_intel.cpp index e290ab90616b94df0135264dcbbf66fe8546a36a..6b7b2c81ebd75bace530cbde5dd9bee1f0981c57 100644 --- a/src/USER-INTEL/dihedral_opls_intel.cpp +++ b/src/USER-INTEL/dihedral_opls_intel.cpp @@ -21,6 +21,7 @@ #include "atom.h" #include "comm.h" #include "memory.h" +#include "modify.h" #include "neighbor.h" #include "domain.h" #include "force.h" diff --git a/src/USER-INTEL/fix_intel.cpp b/src/USER-INTEL/fix_intel.cpp index eac48b8510b4aa67a5e93f689ae1381510e02b57..3e36c8f7a937b0e254389c4099897a62aa9ac236 100644 --- a/src/USER-INTEL/fix_intel.cpp +++ b/src/USER-INTEL/fix_intel.cpp @@ -32,6 +32,7 @@ #include <string.h> #include <stdlib.h> #include <stdio.h> +#include <math.h> #ifdef _LMP_INTEL_OFFLOAD #ifndef INTEL_OFFLOAD_NOAFFINITY diff --git a/src/USER-INTEL/nbin_intel.cpp b/src/USER-INTEL/nbin_intel.cpp index 3a36ead499ebe34a1dfcce5350abbd1f13fb8d6b..9a1dae36ca465b08d8cbbf63407597306b810553 100644 --- a/src/USER-INTEL/nbin_intel.cpp +++ b/src/USER-INTEL/nbin_intel.cpp @@ -18,8 +18,9 @@ #include "nbin_intel.h" #include "atom.h" #include "group.h" -#include "domain.h" #include "comm.h" +#include "domain.h" +#include "modify.h" #include "update.h" #include "error.h" diff --git a/src/USER-INTEL/npair_full_bin_ghost_intel.cpp b/src/USER-INTEL/npair_full_bin_ghost_intel.cpp index e6d45d7b2c6941fa87c24cd089ee46863b4828ee..a814891f2599ad73fcac0d8db7064e56ce173168 100644 --- a/src/USER-INTEL/npair_full_bin_ghost_intel.cpp +++ b/src/USER-INTEL/npair_full_bin_ghost_intel.cpp @@ -21,6 +21,8 @@ #include "neigh_list.h" #include "atom.h" #include "atom_vec.h" +#include "comm.h" +#include "domain.h" #include "molecule.h" #include "error.h" diff --git a/src/USER-INTEL/npair_full_bin_intel.cpp b/src/USER-INTEL/npair_full_bin_intel.cpp index 06c10c080fd9e506760c0d1e66c93ff76233ee7c..60b912d796146d2c9a7f7a15d116afdc430c20ff 100644 --- a/src/USER-INTEL/npair_full_bin_intel.cpp +++ b/src/USER-INTEL/npair_full_bin_intel.cpp @@ -20,6 +20,7 @@ #include "neigh_list.h" #include "atom.h" #include "comm.h" +#include "domain.h" #include "group.h" using namespace LAMMPS_NS; diff --git a/src/USER-INTEL/npair_half_bin_newton_intel.cpp b/src/USER-INTEL/npair_half_bin_newton_intel.cpp index c761557097cf96b677dc20a355a94e186ac5d8c5..8c024a46046827985db3f8047446ad06df61eba1 100644 --- a/src/USER-INTEL/npair_half_bin_newton_intel.cpp +++ b/src/USER-INTEL/npair_half_bin_newton_intel.cpp @@ -20,6 +20,7 @@ #include "neigh_list.h" #include "atom.h" #include "comm.h" +#include "domain.h" #include "group.h" using namespace LAMMPS_NS; diff --git a/src/USER-INTEL/npair_half_bin_newton_tri_intel.cpp b/src/USER-INTEL/npair_half_bin_newton_tri_intel.cpp index d70f1ec5897a0d7b1e8a7b5b9633f9882ce2c54f..653a95139498dbe1d34b5c8d77ca706664e02c56 100644 --- a/src/USER-INTEL/npair_half_bin_newton_tri_intel.cpp +++ b/src/USER-INTEL/npair_half_bin_newton_tri_intel.cpp @@ -20,6 +20,7 @@ #include "neigh_list.h" #include "atom.h" #include "comm.h" +#include "domain.h" #include "group.h" using namespace LAMMPS_NS; diff --git a/src/USER-INTEL/npair_intel.cpp b/src/USER-INTEL/npair_intel.cpp index 0068e02635c7d6d7bd145fe45a35bade50f64f0c..234557c941c1ce7d765e745fb918d6a42297c033 100644 --- a/src/USER-INTEL/npair_intel.cpp +++ b/src/USER-INTEL/npair_intel.cpp @@ -15,6 +15,10 @@ Contributing author: W. Michael Brown (Intel) ------------------------------------------------------------------------- */ +#include "comm.h" +#include "domain.h" +#include "timer.h" +#include "modify.h" #include "npair_intel.h" #include "nstencil.h" diff --git a/src/USER-INTEL/pair_eam_intel.cpp b/src/USER-INTEL/pair_eam_intel.cpp index b97128bf9ff434fbc71c02d204c25128e9583a20..ea4ee30d52602440b06b81603a3d93897b72cba3 100644 --- a/src/USER-INTEL/pair_eam_intel.cpp +++ b/src/USER-INTEL/pair_eam_intel.cpp @@ -23,6 +23,7 @@ #include "atom.h" #include "force.h" #include "comm.h" +#include "modify.h" #include "neighbor.h" #include "neigh_list.h" #include "neigh_request.h" diff --git a/src/USER-INTEL/pppm_disp_intel.cpp b/src/USER-INTEL/pppm_disp_intel.cpp index 1269579ff4f26addaa08d00c19ff433626d5a089..bd41f8b531cc18be1e6cf79376de773a52c9d7c3 100644 --- a/src/USER-INTEL/pppm_disp_intel.cpp +++ b/src/USER-INTEL/pppm_disp_intel.cpp @@ -20,7 +20,10 @@ #include <math.h> #include "pppm_disp_intel.h" #include "atom.h" +#include "comm.h" +#include "domain.h" #include "error.h" +#include "modify.h" #include "fft3d_wrap.h" #include "gridcomm.h" #include "math_const.h" diff --git a/src/USER-INTEL/pppm_intel.cpp b/src/USER-INTEL/pppm_intel.cpp index db855b75ef8f877303ae6186e12fc59314729397..30f8f4c5c0677890115a5f5668a557f2053bb054 100644 --- a/src/USER-INTEL/pppm_intel.cpp +++ b/src/USER-INTEL/pppm_intel.cpp @@ -23,7 +23,10 @@ #include <math.h> #include "pppm_intel.h" #include "atom.h" +#include "comm.h" +#include "domain.h" #include "error.h" +#include "modify.h" #include "fft3d_wrap.h" #include "gridcomm.h" #include "math_const.h" diff --git a/src/USER-OMP/pair_lj_cut_thole_long_omp.cpp b/src/USER-OMP/pair_lj_cut_thole_long_omp.cpp index 85c7e44f8c8654d4eb9e7128c93f8a5e964c309c..110b8917cf07366b02aa0eeda1b99daeaffbe99a 100644 --- a/src/USER-OMP/pair_lj_cut_thole_long_omp.cpp +++ b/src/USER-OMP/pair_lj_cut_thole_long_omp.cpp @@ -29,6 +29,7 @@ #include "math_const.h" #include "error.h" #include "suffix.h" +#include "domain.h" using namespace LAMMPS_NS; using namespace MathConst; diff --git a/src/USER-TALLY/compute_force_tally.cpp b/src/USER-TALLY/compute_force_tally.cpp index 5f29aea5b22cd3b892596ec4f3356e9184a85602..3ec6c188fc385026ce3fd93383fa87e38559b15f 100644 --- a/src/USER-TALLY/compute_force_tally.cpp +++ b/src/USER-TALLY/compute_force_tally.cpp @@ -12,6 +12,7 @@ ------------------------------------------------------------------------- */ #include <string.h> +#include <math.h> #include "compute_force_tally.h" #include "atom.h" #include "group.h" @@ -20,6 +21,7 @@ #include "memory.h" #include "error.h" #include "force.h" +#include "comm.h" using namespace LAMMPS_NS; diff --git a/src/USER-TALLY/compute_heat_flux_tally.cpp b/src/USER-TALLY/compute_heat_flux_tally.cpp index c090050b1540925d3ca76cb3e4f23fcaf7471c4e..43b663b27a580495943b5d4835650ae01811b0cd 100644 --- a/src/USER-TALLY/compute_heat_flux_tally.cpp +++ b/src/USER-TALLY/compute_heat_flux_tally.cpp @@ -20,6 +20,7 @@ #include "memory.h" #include "error.h" #include "force.h" +#include "comm.h" using namespace LAMMPS_NS; diff --git a/src/USER-TALLY/compute_pe_mol_tally.cpp b/src/USER-TALLY/compute_pe_mol_tally.cpp index 25a172b7f81eb6cd68a75fcb520c401bb89ec1de..0328740e0395bc6ca7914ee8d7b7b88e212e0000 100644 --- a/src/USER-TALLY/compute_pe_mol_tally.cpp +++ b/src/USER-TALLY/compute_pe_mol_tally.cpp @@ -20,6 +20,7 @@ #include "memory.h" #include "error.h" #include "force.h" +#include "comm.h" using namespace LAMMPS_NS; diff --git a/src/USER-TALLY/compute_pe_tally.cpp b/src/USER-TALLY/compute_pe_tally.cpp index 5b4644d4e16833f8177a0d22aa29de9802877bc1..caa4cf134a89dd10d73169eae5ba02510a8bf2dd 100644 --- a/src/USER-TALLY/compute_pe_tally.cpp +++ b/src/USER-TALLY/compute_pe_tally.cpp @@ -20,6 +20,7 @@ #include "memory.h" #include "error.h" #include "force.h" +#include "comm.h" using namespace LAMMPS_NS; diff --git a/src/USER-TALLY/compute_stress_tally.cpp b/src/USER-TALLY/compute_stress_tally.cpp index 32253d2cad78857a31c52753f85b856e7b84cec2..e44313d695c46959b87ddb89bbf50ec99078bd7c 100644 --- a/src/USER-TALLY/compute_stress_tally.cpp +++ b/src/USER-TALLY/compute_stress_tally.cpp @@ -20,6 +20,8 @@ #include "memory.h" #include "error.h" #include "force.h" +#include "comm.h" +#include "domain.h" using namespace LAMMPS_NS; diff --git a/src/accelerator_kokkos.h b/src/accelerator_kokkos.h index 8ea5b9d4d2dd93dd4f85147035bfc5be1633024a..708592a25eb000cf3bdc5bb38add4b413809a94b 100644 --- a/src/accelerator_kokkos.h +++ b/src/accelerator_kokkos.h @@ -25,6 +25,7 @@ #include "comm_tiled_kokkos.h" #include "domain_kokkos.h" #include "neighbor_kokkos.h" +#include "memory_kokkos.h" #include "modify_kokkos.h" #else @@ -37,6 +38,7 @@ #include "comm_tiled.h" #include "domain.h" #include "neighbor.h" +#include "memory.h" #include "modify.h" namespace LAMMPS_NS { @@ -89,6 +91,13 @@ class NeighborKokkos : public Neighbor { ~NeighborKokkos() {} }; +class MemoryKokkos : public Memory { + public: + MemoryKokkos(class LAMMPS *lmp) : Memory(lmp) {} + ~MemoryKokkos() {} + void grow_kokkos(tagint **, tagint **, int, int, const char*) {} +}; + class ModifyKokkos : public Modify { public: ModifyKokkos(class LAMMPS *lmp) : Modify(lmp) {} diff --git a/src/compute_aggregate_atom.cpp b/src/compute_aggregate_atom.cpp index 1155ac437adfb439cda2813ea0b8d4f4a850afed..1e91327e54639be94bf89810058c7d9e1df5e2a9 100644 --- a/src/compute_aggregate_atom.cpp +++ b/src/compute_aggregate_atom.cpp @@ -16,6 +16,7 @@ ------------------------------------------------------------------------- */ #include <string.h> +#include <math.h> #include "compute_aggregate_atom.h" #include "atom.h" #include "atom_vec.h" diff --git a/src/compute_orientorder_atom.cpp b/src/compute_orientorder_atom.cpp index 90e2830e39e55881884bb6cb57f137ec387f62bf..b443d56bf82b3fd7a8d4f50b58181ecb49e922e4 100644 --- a/src/compute_orientorder_atom.cpp +++ b/src/compute_orientorder_atom.cpp @@ -18,6 +18,7 @@ #include <string.h> #include <stdlib.h> +#include <math.h> #include "compute_orientorder_atom.h" #include "atom.h" #include "update.h" diff --git a/src/compute_rdf.cpp b/src/compute_rdf.cpp index 167de4576dfcbe7c24d8db5c1049c84806c29ecb..bcb620f3b36c59dd842669dd626baebcda760f31 100644 --- a/src/compute_rdf.cpp +++ b/src/compute_rdf.cpp @@ -32,6 +32,7 @@ #include "math_const.h" #include "memory.h" #include "error.h" +#include "comm.h" using namespace LAMMPS_NS; using namespace MathConst; 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/fix_addforce.cpp b/src/fix_addforce.cpp index 5c677a43385a071a99d1e89489b6cb2a6ace6e8c..6b1e89227985acb0dc12336520962ff7599ffee3 100644 --- a/src/fix_addforce.cpp +++ b/src/fix_addforce.cpp @@ -16,7 +16,6 @@ #include "fix_addforce.h" #include "atom.h" #include "atom_masks.h" -#include "accelerator_kokkos.h" #include "update.h" #include "modify.h" #include "domain.h" diff --git a/src/input.cpp b/src/input.cpp index d18be498fd50d062866cf5ec2f50f14cb1341a9c..06403276824a471aaafdc3453706ca951624c657 100644 --- a/src/input.cpp +++ b/src/input.cpp @@ -27,7 +27,6 @@ #include "comm.h" #include "comm_brick.h" #include "comm_tiled.h" -#include "accelerator_kokkos.h" #include "group.h" #include "domain.h" #include "output.h" 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/lammps.h b/src/lammps.h index c432784a0b14c92f571dc032040f22e35ec30f96..e3b815e6237a7ad660b90b97ee7c65b64db858db 100644 --- a/src/lammps.h +++ b/src/lammps.h @@ -53,6 +53,7 @@ class LAMMPS { class KokkosLMP *kokkos; // KOKKOS accelerator class class AtomKokkos *atomKK; // KOKKOS version of Atom class + class MemoryKokkos *memoryKK; // KOKKOS version of Memory class class Python * python; // Python interface diff --git a/src/memory.h b/src/memory.h index b83482e4c2b69c9d11259e7c5d1ea73ba7ba077b..f2faecf6e177587e4aaf9dfc14991c6ef4f94f58 100644 --- a/src/memory.h +++ b/src/memory.h @@ -16,9 +16,6 @@ #include "lmptype.h" #include "pointers.h" -#ifdef LMP_KOKKOS -#include "kokkos_type.h" -#endif namespace LAMMPS_NS { @@ -31,16 +28,6 @@ class Memory : protected Pointers { void sfree(void *); void fail(const char *); - // Kokkos memory allocation functions - // provide a dummy prototpye for any Kokkos memory function - // called in main LAMMPS even when not built with KOKKOS package - -#ifdef LMP_KOKKOS -#include "memory_kokkos.h" -#else - void grow_kokkos(tagint **, tagint **, int, int, const char*) {} -#endif - /* ---------------------------------------------------------------------- create/grow/destroy vecs and multidim arrays with contiguous memory blocks only use with primitive data types, e.g. 1d vec of ints, 2d array of doubles diff --git a/src/pair.h b/src/pair.h index eb71e8822474aed87e024988e27b669196385cd9..cfb6576653c3aca4d21dfa99a0a6feda78a4977a 100644 --- a/src/pair.h +++ b/src/pair.h @@ -15,7 +15,6 @@ #define LMP_PAIR_H #include "pointers.h" -#include "accelerator_kokkos.h" namespace LAMMPS_NS { @@ -165,10 +164,6 @@ class Pair : protected Pointers { virtual int pack_forward_comm(int, int *, double *, int, int *) {return 0;} virtual void unpack_forward_comm(int, int, double *) {} - 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 &) {} virtual int pack_reverse_comm(int, int, double *) {return 0;} virtual void unpack_reverse_comm(int, int *, double *) {} virtual double memory_usage(); diff --git a/src/pointers.h b/src/pointers.h index 82b49c1dad05e6988ec02bd1b5646ee2078b6442..44967f51359a28067fee329362a65c8b3a8920dd 100644 --- a/src/pointers.h +++ b/src/pointers.h @@ -57,6 +57,7 @@ class Pointers { screen(ptr->screen), logfile(ptr->logfile), atomKK(ptr->atomKK), + memoryKK(ptr->memoryKK), python(ptr->python) {} virtual ~Pointers() {} @@ -84,6 +85,7 @@ class Pointers { FILE *&logfile; class AtomKokkos *&atomKK; + class MemoryKokkos *&memoryKK; class Python *&python; }; 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; diff --git a/src/special.cpp b/src/special.cpp index 381a763dd29ad9cec5254e8556763f0967d826ac..56529e748b8d52e070696543627bfba255b87e0a 100644 --- a/src/special.cpp +++ b/src/special.cpp @@ -591,7 +591,8 @@ void Special::combine() AtomKokkos* atomKK = (AtomKokkos*) atom; atomKK->modified(Host,SPECIAL_MASK); atomKK->sync(Device,SPECIAL_MASK); - memory->grow_kokkos(atomKK->k_special,atom->special, + MemoryKokkos* memoryKK = (MemoryKokkos*) memory; + memoryKK->grow_kokkos(atomKK->k_special,atom->special, atom->nmax,atom->maxspecial,"atom:special"); atomKK->modified(Device,SPECIAL_MASK); atomKK->sync(Host,SPECIAL_MASK); diff --git a/src/write_coeff.cpp b/src/write_coeff.cpp index d4d4f99bf388c597f9ef5b674b06ec07a933871d..0556d647cf36d1046f4864002f2e6fa5e5d656f5 100644 --- a/src/write_coeff.cpp +++ b/src/write_coeff.cpp @@ -23,6 +23,7 @@ #include "force.h" #include "universe.h" #include "error.h" +#include "domain.h" using namespace LAMMPS_NS;