diff --git a/src/KOKKOS/Install.sh b/src/KOKKOS/Install.sh index bceb82a0b3..102e075124 100644 --- a/src/KOKKOS/Install.sh +++ b/src/KOKKOS/Install.sh @@ -29,6 +29,10 @@ touch ../memory.h # list of files with optional dependcies +action angle_charmm_kokkos.cpp angle_charmm.cpp +action angle_charmm_kokkos.h angle_charmm.h +action angle_harmonic_kokkos.cpp angle_harmonic.cpp +action angle_harmonic_kokkos.h angle_harmonic.h action atom_kokkos.cpp action atom_kokkos.h action atom_vec_angle_kokkos.cpp atom_vec_angle.cpp @@ -45,20 +49,32 @@ action atom_vec_kokkos.cpp action atom_vec_kokkos.h action atom_vec_molecular_kokkos.cpp atom_vec_molecular.cpp action atom_vec_molecular_kokkos.h atom_vec_molecular.h +action bond_fene_kokkos.cpp bond_fene.cpp +action bond_fene_kokkos.h bond_fene.h +action bond_harmonic_kokkos.cpp bond_harmonic.cpp +action bond_harmonic_kokkos.h bond_harmonic.h action comm_kokkos.cpp action comm_kokkos.h +action dihedral_charmm_kokkos.cpp dihedral_charmm.cpp +action dihedral_charmm_kokkos.h dihedral_charmm.h +action dihedral_opls_kokkos.cpp dihedral_opls.cpp +action dihedral_opls_kokkos.h dihedral_opls.h action domain_kokkos.cpp action domain_kokkos.h action fix_langevin_kokkos.cpp action fix_langevin_kokkos.h action fix_nve_kokkos.cpp action fix_nve_kokkos.h +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 memory_kokkos.h action modify_kokkos.cpp action modify_kokkos.h +action neigh_bond_kokkos.cpp +action neigh_bond_kokkos.h action neigh_full_kokkos.h action neigh_list_kokkos.cpp action neigh_list_kokkos.h diff --git a/src/KOKKOS/atom_vec_angle_kokkos.cpp b/src/KOKKOS/atom_vec_angle_kokkos.cpp index e1f73064d7..d78014120f 100644 --- a/src/KOKKOS/atom_vec_angle_kokkos.cpp +++ b/src/KOKKOS/atom_vec_angle_kokkos.cpp @@ -1636,6 +1636,7 @@ void AtomVecAngleKokkos::data_atom(double *coord, imageint imagetmp, { int nlocal = atom->nlocal; if (nlocal == nmax) grow(0); + atomKK->modified(Host,ALL_MASK); h_tag(nlocal) = atoi(values[0]); h_molecule(nlocal) = atoi(values[1]); diff --git a/src/KOKKOS/atom_vec_bond_kokkos.cpp b/src/KOKKOS/atom_vec_bond_kokkos.cpp index a5ed6163a7..a991c594f3 100644 --- a/src/KOKKOS/atom_vec_bond_kokkos.cpp +++ b/src/KOKKOS/atom_vec_bond_kokkos.cpp @@ -1507,6 +1507,7 @@ void AtomVecBondKokkos::data_atom(double *coord, imageint imagetmp, { int nlocal = atomKK->nlocal; if (nlocal == nmax) grow(0); + atomKK->modified(Host,ALL_MASK); h_tag(nlocal) = atoi(values[0]); h_molecule(nlocal) = atoi(values[1]); diff --git a/src/KOKKOS/atom_vec_full_kokkos.cpp b/src/KOKKOS/atom_vec_full_kokkos.cpp index 6623fa2f25..0d3893432e 100644 --- a/src/KOKKOS/atom_vec_full_kokkos.cpp +++ b/src/KOKKOS/atom_vec_full_kokkos.cpp @@ -1970,6 +1970,7 @@ void AtomVecFullKokkos::data_atom(double *coord, imageint imagetmp, { int nlocal = atom->nlocal; if (nlocal == nmax) grow(0); + atomKK->modified(Host,ALL_MASK); h_tag(nlocal) = atoi(values[0]); h_molecule(nlocal) = atoi(values[1]); @@ -2176,7 +2177,7 @@ void AtomVecFullKokkos::sync(ExecutionSpace space, unsigned int mask) atomKK->k_improper_atom1.sync(); atomKK->k_improper_atom2.sync(); atomKK->k_improper_atom3.sync(); - atomKK->k_improper_atom3.sync(); + atomKK->k_improper_atom4.sync(); } } else { if (mask & X_MASK) atomKK->k_x.sync(); @@ -2218,7 +2219,7 @@ void AtomVecFullKokkos::sync(ExecutionSpace space, unsigned int mask) atomKK->k_improper_atom1.sync(); atomKK->k_improper_atom2.sync(); atomKK->k_improper_atom3.sync(); - atomKK->k_improper_atom3.sync(); + atomKK->k_improper_atom4.sync(); } } } @@ -2267,7 +2268,7 @@ void AtomVecFullKokkos::modified(ExecutionSpace space, unsigned int mask) atomKK->k_improper_atom1.modify(); atomKK->k_improper_atom2.modify(); atomKK->k_improper_atom3.modify(); - atomKK->k_improper_atom3.modify(); + atomKK->k_improper_atom4.modify(); } } else { if (mask & X_MASK) atomKK->k_x.modify(); @@ -2309,7 +2310,7 @@ void AtomVecFullKokkos::modified(ExecutionSpace space, unsigned int mask) atomKK->k_improper_atom1.modify(); atomKK->k_improper_atom2.modify(); atomKK->k_improper_atom3.modify(); - atomKK->k_improper_atom3.modify(); + atomKK->k_improper_atom4.modify(); } } } diff --git a/src/KOKKOS/atom_vec_molecular_kokkos.cpp b/src/KOKKOS/atom_vec_molecular_kokkos.cpp index 7c48b2dc85..256514eb7f 100644 --- a/src/KOKKOS/atom_vec_molecular_kokkos.cpp +++ b/src/KOKKOS/atom_vec_molecular_kokkos.cpp @@ -1895,6 +1895,7 @@ void AtomVecMolecularKokkos::data_atom(double *coord, imageint imagetmp, { int nlocal = atom->nlocal; if (nlocal == nmax) grow(0); + atomKK->modified(Host,ALL_MASK); h_tag(nlocal) = atoi(values[0]); h_molecule(nlocal) = atoi(values[1]); @@ -2094,7 +2095,7 @@ void AtomVecMolecularKokkos::sync(ExecutionSpace space, unsigned int mask) atomKK->k_improper_atom1.sync(); atomKK->k_improper_atom2.sync(); atomKK->k_improper_atom3.sync(); - atomKK->k_improper_atom3.sync(); + atomKK->k_improper_atom4.sync(); } } else { if (mask & X_MASK) atomKK->k_x.sync(); @@ -2135,7 +2136,7 @@ void AtomVecMolecularKokkos::sync(ExecutionSpace space, unsigned int mask) atomKK->k_improper_atom1.sync(); atomKK->k_improper_atom2.sync(); atomKK->k_improper_atom3.sync(); - atomKK->k_improper_atom3.sync(); + atomKK->k_improper_atom4.sync(); } } } @@ -2183,7 +2184,7 @@ void AtomVecMolecularKokkos::modified(ExecutionSpace space, unsigned int mask) atomKK->k_improper_atom1.modify(); atomKK->k_improper_atom2.modify(); atomKK->k_improper_atom3.modify(); - atomKK->k_improper_atom3.modify(); + atomKK->k_improper_atom4.modify(); } } else { if (mask & X_MASK) atomKK->k_x.modify(); @@ -2224,7 +2225,7 @@ void AtomVecMolecularKokkos::modified(ExecutionSpace space, unsigned int mask) atomKK->k_improper_atom1.modify(); atomKK->k_improper_atom2.modify(); atomKK->k_improper_atom3.modify(); - atomKK->k_improper_atom3.modify(); + atomKK->k_improper_atom4.modify(); } } } diff --git a/src/KOKKOS/comm_kokkos.cpp b/src/KOKKOS/comm_kokkos.cpp index 8c3b825acd..bcbf1efc7e 100644 --- a/src/KOKKOS/comm_kokkos.cpp +++ b/src/KOKKOS/comm_kokkos.cpp @@ -888,10 +888,11 @@ void CommKokkos::borders_device() { // reset global->local map - if (map_style) atom->map_set(); if (exec_space == Host) k_sendlist.sync(); atomKK->modified(exec_space,ALL_MASK); DeviceType::fence(); + atomKK->sync(Host,TAG_MASK); + if (map_style) atom->map_set(); } /* ---------------------------------------------------------------------- realloc the size of the send buffer as needed with BUFFACTOR and bufextra diff --git a/src/KOKKOS/kokkos_type.h b/src/KOKKOS/kokkos_type.h index 4f4f5f45c3..ebedfbadcc 100644 --- a/src/KOKKOS/kokkos_type.h +++ b/src/KOKKOS/kokkos_type.h @@ -232,7 +232,7 @@ typedef tdual_tagint_1d::t_dev_const_um t_tagint_1d_const_um; typedef tdual_tagint_1d::t_dev_const_randomread t_tagint_1d_randomread; typedef Kokkos:: - DualView + DualView tdual_tagint_2d; typedef tdual_tagint_2d::t_dev t_tagint_2d; typedef tdual_tagint_2d::t_dev_const t_tagint_2d_const; @@ -455,7 +455,7 @@ typedef tdual_tagint_1d::t_host_const_um t_tagint_1d_const_um; typedef tdual_tagint_1d::t_host_const_randomread t_tagint_1d_randomread; typedef Kokkos:: - DualView + DualView tdual_tagint_2d; typedef tdual_tagint_2d::t_host t_tagint_2d; typedef tdual_tagint_2d::t_host_const t_tagint_2d_const; diff --git a/src/KOKKOS/neighbor_kokkos.cpp b/src/KOKKOS/neighbor_kokkos.cpp index 133ac7cea0..493a5e546e 100644 --- a/src/KOKKOS/neighbor_kokkos.cpp +++ b/src/KOKKOS/neighbor_kokkos.cpp @@ -16,6 +16,10 @@ #include "pair.h" #include "neigh_request.h" #include "memory.h" +#include "update.h" +#include "atom_masks.h" +#include "error.h" +#include "kokkos.h" using namespace LAMMPS_NS; @@ -23,7 +27,8 @@ enum{NSQ,BIN,MULTI}; // also in neigh_list.cpp /* ---------------------------------------------------------------------- */ -NeighborKokkos::NeighborKokkos(LAMMPS *lmp) : Neighbor(lmp) +NeighborKokkos::NeighborKokkos(LAMMPS *lmp) : Neighbor(lmp), + neighbond_device(lmp),neighbond_host(lmp) { atoms_per_bin = 16; @@ -35,32 +40,40 @@ NeighborKokkos::NeighborKokkos(LAMMPS *lmp) : Neighbor(lmp) lists_device = NULL; pair_build_device = NULL; stencil_create_device = NULL; + + device_flag = 0; } /* ---------------------------------------------------------------------- */ NeighborKokkos::~NeighborKokkos() { - memory->destroy_kokkos(k_cutneighsq,cutneighsq); - cutneighsq = NULL; - - for (int i = 0; i < nlist_host; i++) delete lists_host[i]; - delete [] lists_host; - for (int i = 0; i < nlist_device; i++) delete lists_device[i]; - delete [] lists_device; - - delete [] pair_build_device; - delete [] pair_build_host; - - 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); - + if (!copymode) { + memory->destroy_kokkos(k_cutneighsq,cutneighsq); + cutneighsq = NULL; + + for (int i = 0; i < nlist_host; i++) delete lists_host[i]; + delete [] lists_host; + for (int i = 0; i < nlist_device; i++) delete lists_device[i]; + delete [] lists_device; + + delete [] pair_build_device; + delete [] pair_build_host; + + 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_bondlist,bondlist); + memory->destroy_kokkos(k_anglelist,anglelist); + memory->destroy_kokkos(k_dihedrallist,dihedrallist); + memory->destroy_kokkos(k_improperlist,improperlist); + } } /* ---------------------------------------------------------------------- */ @@ -144,6 +157,11 @@ int NeighborKokkos::init_lists_kokkos() } } + // 1st time allocation of xhold + + if (dist_check) + xhold = DAT::tdual_x_array("neigh:xhold",maxhold); + // return # of non-Kokkos lists return nlist; @@ -270,14 +288,204 @@ void NeighborKokkos::choose_build(int index, NeighRequest *rq) Neighbor::choose_build(index,rq); } -/* ---------------------------------------------------------------------- */ +/* ---------------------------------------------------------------------- + if any atom moved trigger distance (half of neighbor skin) return 1 + shrink trigger distance if box size has changed + conservative shrink procedure: + compute distance each of 8 corners of box has moved since last reneighbor + reduce skin distance by sum of 2 largest of the 8 values + new trigger = 1/2 of reduced skin distance + for orthogonal box, only need 2 lo/hi corners + for triclinic, need all 8 corners since deformations can displace all 8 +------------------------------------------------------------------------- */ -void NeighborKokkos::build_kokkos(int i) +int NeighborKokkos::check_distance() { - if (lists_host[blist[i]]) - (this->*pair_build_host[blist[i]])(lists_host[blist[i]]); - else if (lists_device[blist[i]]) - (this->*pair_build_device[blist[i]])(lists_device[blist[i]]); + if (nlist_device) + check_distance_kokkos(); + else + check_distance_kokkos(); +} + +template +int NeighborKokkos::check_distance_kokkos() +{ + typedef DeviceType device_type; + + double delx,dely,delz,rsq; + double delta,delta1,delta2; + + if (boxcheck) { + if (triclinic == 0) { + delx = bboxlo[0] - boxlo_hold[0]; + dely = bboxlo[1] - boxlo_hold[1]; + delz = bboxlo[2] - boxlo_hold[2]; + delta1 = sqrt(delx*delx + dely*dely + delz*delz); + delx = bboxhi[0] - boxhi_hold[0]; + dely = bboxhi[1] - boxhi_hold[1]; + delz = bboxhi[2] - boxhi_hold[2]; + delta2 = sqrt(delx*delx + dely*dely + delz*delz); + delta = 0.5 * (skin - (delta1+delta2)); + deltasq = delta*delta; + } else { + domain->box_corners(); + delta1 = delta2 = 0.0; + for (int i = 0; i < 8; i++) { + delx = corners[i][0] - corners_hold[i][0]; + dely = corners[i][1] - corners_hold[i][1]; + delz = corners[i][2] - corners_hold[i][2]; + delta = sqrt(delx*delx + dely*dely + delz*delz); + if (delta > delta1) delta1 = delta; + else if (delta > delta2) delta2 = delta; + } + delta = 0.5 * (skin - (delta1+delta2)); + deltasq = delta*delta; + } + } else deltasq = triggersq; + + atomKK->sync(ExecutionSpaceFromDevice::space,X_MASK); + x = atomKK->k_x; + xhold.sync(); + int nlocal = atom->nlocal; + if (includegroup) nlocal = atom->nfirst; + + int flag = 0; + copymode = 1; + Kokkos::parallel_reduce(Kokkos::RangePolicy >(0,nlocal),*this,flag); + DeviceType::fence(); + copymode = 0; + + int flagall; + MPI_Allreduce(&flag,&flagall,1,MPI_INT,MPI_MAX,world); + if (flagall && ago == MAX(every,delay)) ndanger++; + return flagall; +} + +template +KOKKOS_INLINE_FUNCTION +void NeighborKokkos::operator()(TagNeighborCheckDistance, const int &i, int &flag) const { + typedef DeviceType device_type; + const X_FLOAT delx = x.view()(i,0) - xhold.view()(i,0); + const X_FLOAT dely = x.view()(i,1) - xhold.view()(i,1); + const X_FLOAT delz = x.view()(i,2) - xhold.view()(i,2); + const X_FLOAT rsq = delx*delx + dely*dely + delz*delz; + if (rsq > deltasq) flag = 1; +} + +/* ---------------------------------------------------------------------- + build perpetuals neighbor lists + called at setup and every few timesteps during run or minimization + topology lists also built if topoflag = 1, USER-CUDA calls with topoflag = 0 +------------------------------------------------------------------------- */ + + +void NeighborKokkos::build(int topoflag) +{ + if (nlist_device) + this->template build_kokkos(topoflag); + else + this->template build_kokkos(topoflag); +} + +template +void NeighborKokkos::build_kokkos(int topoflag) +{ + typedef DeviceType device_type; + + int i; + + ago = 0; + ncalls++; + lastcall = update->ntimestep; + + // store current atom positions and box size if needed + + if (dist_check) { + atomKK->sync(ExecutionSpaceFromDevice::space,X_MASK); + x = atomKK->k_x; + int nlocal = atom->nlocal; + if (includegroup) nlocal = atom->nfirst; + if (nlocal > maxhold) { + maxhold = atom->nmax; + xhold = DAT::tdual_x_array("neigh:xhold",maxhold); + } + copymode = 1; + Kokkos::parallel_for(Kokkos::RangePolicy >(0,nlocal),*this); + DeviceType::fence(); + copymode = 0; + xhold.modify(); + if (boxcheck) { + if (triclinic == 0) { + boxlo_hold[0] = bboxlo[0]; + boxlo_hold[1] = bboxlo[1]; + boxlo_hold[2] = bboxlo[2]; + boxhi_hold[0] = bboxhi[0]; + boxhi_hold[1] = bboxhi[1]; + boxhi_hold[2] = bboxhi[2]; + } else { + domain->box_corners(); + corners = domain->corners; + for (i = 0; i < 8; i++) { + corners_hold[i][0] = corners[i][0]; + corners_hold[i][1] = corners[i][1]; + corners_hold[i][2] = corners[i][2]; + } + } + } + } + + // if any lists store neighbors of ghosts: + // invoke grow() if nlocal+nghost exceeds previous list size + // else only invoke grow() if nlocal exceeds previous list size + // only for lists with growflag set and which are perpetual (glist) + + if (anyghostlist && atom->nlocal+atom->nghost > maxatom) { + maxatom = atom->nmax; + for (i = 0; i < nglist; i++) lists[glist[i]]->grow(maxatom); + } else if (atom->nlocal > maxatom) { + maxatom = atom->nmax; + for (i = 0; i < nglist; i++) lists[glist[i]]->grow(maxatom); + } + + // extend atom bin list if necessary + + if (style != NSQ && atom->nmax > maxbin) { + maxbin = atom->nmax; + memory->destroy(bins); + memory->create(bins,maxbin,"bins"); + } + + // check that using special bond flags will not overflow neigh lists + + if (atom->nlocal+atom->nghost > NEIGHMASK) + error->one(FLERR,"Too many local+ghost atoms for neighbor list"); + + // invoke building of pair and molecular topology neighbor lists + // only for pairwise lists with buildflag set + // blist is for standard neigh lists, otherwise is a Kokkos list + + for (i = 0; i < nblist; i++) { + if (lists[blist[i]]) + (this->*pair_build[blist[i]])(lists[blist[i]]); + else { + if (lists_host[blist[i]]) + (this->*pair_build_host[blist[i]])(lists_host[blist[i]]); + else if (lists_device[blist[i]]) + (this->*pair_build_device[blist[i]])(lists_device[blist[i]]); + } + } + + if (atom->molecular && topoflag) + build_topology_kokkos(); +} + +template +KOKKOS_INLINE_FUNCTION +void NeighborKokkos::operator()(TagNeighborXhold, const int &i) const { + typedef DeviceType device_type; + xhold.view()(i,0) = x.view()(i,0); + xhold.view()(i,1) = x.view()(i,1); + xhold.view()(i,2) = x.view()(i,2); } /* ---------------------------------------------------------------------- */ @@ -323,6 +531,49 @@ void NeighborKokkos::modify_mol_group_grow_kokkos(){ k_ex_mol_group.modify(); } +/* ---------------------------------------------------------------------- */ + +void NeighborKokkos::init_topology_kokkos() { + if (nlist_device) { + neighbond_device.init_topology_kk(); + } else { + neighbond_host.init_topology_kk(); + } +} + +/* ---------------------------------------------------------------------- + build all topology neighbor lists every few timesteps + normally built with pair lists, but USER-CUDA separates them +------------------------------------------------------------------------- */ + +void NeighborKokkos::build_topology_kokkos() { + if (nlist_device) { + neighbond_device.build_topology_kk(); + + k_bondlist = neighbond_device.k_bondlist; + k_anglelist = neighbond_device.k_anglelist; + k_dihedrallist = neighbond_device.k_dihedrallist; + k_improperlist = neighbond_device.k_improperlist; + + k_bondlist.modify(); + k_anglelist.modify(); + k_dihedrallist.modify(); + k_improperlist.modify(); + } else { + neighbond_host.build_topology_kk(); + + k_bondlist = neighbond_host.k_bondlist; + k_anglelist = neighbond_host.k_anglelist; + k_dihedrallist = neighbond_host.k_dihedrallist; + k_improperlist = neighbond_host.k_improperlist; + + k_bondlist.modify(); + k_anglelist.modify(); + k_dihedrallist.modify(); + k_improperlist.modify(); + } +} + // include to trigger instantiation of templated functions #include "neigh_full_kokkos.h" diff --git a/src/KOKKOS/neighbor_kokkos.h b/src/KOKKOS/neighbor_kokkos.h index 03c52dc030..a461c2d0f8 100644 --- a/src/KOKKOS/neighbor_kokkos.h +++ b/src/KOKKOS/neighbor_kokkos.h @@ -16,6 +16,7 @@ #include "neighbor.h" #include "neigh_list_kokkos.h" +#include "neigh_bond_kokkos.h" #include "kokkos_type.h" #include @@ -274,8 +275,16 @@ struct NeighborClusterKokkosBuildFunctor { } }; +template +struct TagNeighborCheckDistance{}; + +template +struct TagNeighborXhold{}; + class NeighborKokkos : public Neighbor { public: + typedef int value_type; + class AtomKokkos *atomKK; int nlist_host; // pairwise neighbor lists on Host @@ -283,10 +292,26 @@ class NeighborKokkos : public Neighbor { int nlist_device; // pairwise neighbor lists on Device NeighListKokkos **lists_device; + NeighBondKokkos neighbond_host; + NeighBondKokkos neighbond_device; + + DAT::tdual_int_2d k_bondlist; + DAT::tdual_int_2d k_anglelist; + DAT::tdual_int_2d k_dihedrallist; + DAT::tdual_int_2d k_improperlist; + NeighborKokkos(class LAMMPS *); ~NeighborKokkos(); void init(); + template + KOKKOS_INLINE_FUNCTION + void operator()(TagNeighborCheckDistance, const int&, int&) const; + + template + KOKKOS_INLINE_FUNCTION + void operator()(TagNeighborXhold, const int&) const; + private: int atoms_per_bin; DAT::tdual_xfloat_2d k_cutneighsq; @@ -300,6 +325,12 @@ class NeighborKokkos : public Neighbor { DAT::tdual_int_1d k_ex_mol_group; DAT::tdual_int_1d k_ex_mol_bit; + DAT::tdual_x_array x; + DAT::tdual_x_array xhold; + + X_FLOAT deltasq; + int device_flag; + void init_cutneighsq_kokkos(int); int init_lists_kokkos(); void init_list_flags1_kokkos(int); @@ -309,11 +340,16 @@ class NeighborKokkos : public Neighbor { void init_ex_bit_kokkos(); void init_ex_mol_bit_kokkos(); void choose_build(int, NeighRequest *); - void build_kokkos(int); + virtual int check_distance(); + template int check_distance_kokkos(); + virtual void build(int); + template void build_kokkos(int); void setup_bins_kokkos(int); void modify_ex_type_grow_kokkos(); void modify_ex_group_grow_kokkos(); void modify_mol_group_grow_kokkos(); + void init_topology_kokkos(); + void build_topology_kokkos(); typedef void (NeighborKokkos::*PairPtrHost) (class NeighListKokkos *);