From f70ac64c5d544d89f30600d36adaee9ab9950fab Mon Sep 17 00:00:00 2001 From: Stan Gerald Moore Date: Fri, 7 Oct 2022 16:25:38 -0600 Subject: [PATCH 1/6] Fix obscure bug in Kokkos neigh list build --- src/KOKKOS/npair_copy_kokkos.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/src/KOKKOS/npair_copy_kokkos.cpp b/src/KOKKOS/npair_copy_kokkos.cpp index 5750b5c50f..55c1d52596 100644 --- a/src/KOKKOS/npair_copy_kokkos.cpp +++ b/src/KOKKOS/npair_copy_kokkos.cpp @@ -63,6 +63,7 @@ void NPairCopyKokkos::copy_to_kokkos(NeighList *list) list_kk->d_ilist = listcopy_kk->d_ilist; list_kk->d_numneigh = listcopy_kk->d_numneigh; list_kk->d_neighbors = listcopy_kk->d_neighbors; + list_kk->maxneighs = listcopy_kk->maxneighs; } /* ---------------------------------------------------------------------- */ From 4ce966e500d70cd79d4e3829b3c4a2338772167c Mon Sep 17 00:00:00 2001 From: Stan Moore Date: Fri, 7 Oct 2022 15:50:31 -0700 Subject: [PATCH 2/6] Prevent segfault from wild pointer --- src/REAXFF/fix_reaxff_species.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/src/REAXFF/fix_reaxff_species.cpp b/src/REAXFF/fix_reaxff_species.cpp index 7307621a83..f8708697f4 100644 --- a/src/REAXFF/fix_reaxff_species.cpp +++ b/src/REAXFF/fix_reaxff_species.cpp @@ -134,6 +134,7 @@ FixReaxFFSpecies::FixReaxFFSpecies(LAMMPS *lmp, int narg, char **arg) : Fix(lmp, nd = nullptr; molmap = nullptr; mark = nullptr; + fdel = nullptr; nmax = 0; setupflag = 0; From e661297838bc6bd89b00b996e069087316b478fa Mon Sep 17 00:00:00 2001 From: Axel Kohlmeyer Date: Sat, 8 Oct 2022 11:33:38 -0400 Subject: [PATCH 3/6] move pointer initialization to initializer list. modernize error messages. --- src/KOKKOS/fix_reaxff_species_kokkos.cpp | 6 +- src/REAXFF/fix_reaxff_species.cpp | 110 +++++++++++------------ src/REAXFF/fix_reaxff_species.h | 2 +- 3 files changed, 57 insertions(+), 61 deletions(-) diff --git a/src/KOKKOS/fix_reaxff_species_kokkos.cpp b/src/KOKKOS/fix_reaxff_species_kokkos.cpp index 2ae14571aa..3bc854d3a0 100644 --- a/src/KOKKOS/fix_reaxff_species_kokkos.cpp +++ b/src/KOKKOS/fix_reaxff_species_kokkos.cpp @@ -54,8 +54,8 @@ FixReaxFFSpeciesKokkos::FixReaxFFSpeciesKokkos(LAMMPS *lmp, int narg, char **arg void FixReaxFFSpeciesKokkos::init() { Pair* pair_kk = force->pair_match("^reax../kk",0); - if (pair_kk == nullptr) error->all(FLERR,"Cannot use fix reaxff/species/kk without " - "pair_style reaxff/kk"); + if (pair_kk == nullptr) + error->all(FLERR,"Cannot use fix reaxff/species/kk without pair_style reaxff/kk"); FixReaxFFSpecies::init(); } @@ -135,7 +135,7 @@ void FixReaxFFSpeciesKokkos::FindMolecule() if (!anychange) break; MPI_Allreduce(&loop,&looptot,1,MPI_INT,MPI_SUM,world); - if (looptot >= 400*nprocs) break; + if (looptot >= 400*comm->nprocs) break; } } diff --git a/src/REAXFF/fix_reaxff_species.cpp b/src/REAXFF/fix_reaxff_species.cpp index f8708697f4..322ef8aadf 100644 --- a/src/REAXFF/fix_reaxff_species.cpp +++ b/src/REAXFF/fix_reaxff_species.cpp @@ -44,9 +44,13 @@ using namespace FixConst; /* ---------------------------------------------------------------------- */ -FixReaxFFSpecies::FixReaxFFSpecies(LAMMPS *lmp, int narg, char **arg) : Fix(lmp, narg, arg) +FixReaxFFSpecies::FixReaxFFSpecies(LAMMPS *lmp, int narg, char **arg) : + Fix(lmp, narg, arg), Name(nullptr), MolName(nullptr), NMol(nullptr), nd(nullptr), + MolType(nullptr), molmap(nullptr), mark(nullptr), Mol2Spec(nullptr), clusterID(nullptr), + x0(nullptr), BOCut(nullptr), fp(nullptr), pos(nullptr), fdel(nullptr), ele(nullptr), + eletype(nullptr), filepos(nullptr), filedel(nullptr) { - if (narg < 7) error->all(FLERR, "Illegal fix reaxff/species command"); + if (narg < 7) utils::missing_cmd_args(FLERR, "fix reaxff/species", error); force_reneighbor = 1; next_reneighbor = -1; @@ -62,8 +66,6 @@ FixReaxFFSpecies::FixReaxFFSpecies(LAMMPS *lmp, int narg, char **arg) : Fix(lmp, compressed = 0; nvalid = -1; - MPI_Comm_rank(world, &me); - MPI_Comm_size(world, &nprocs); ntypes = atom->ntypes; nevery = utils::inumeric(FLERR, arg[3], false, lmp); @@ -73,10 +75,11 @@ FixReaxFFSpecies::FixReaxFFSpecies(LAMMPS *lmp, int narg, char **arg) : Fix(lmp, comm_forward = 4; - if (nevery <= 0 || nrepeat <= 0 || nfreq <= 0) - error->all(FLERR, "Illegal fix reaxff/species command"); - if (nfreq % nevery || nrepeat * nevery > nfreq) - error->all(FLERR, "Illegal fix reaxff/species command"); + if (nevery <= 0) error->all(FLERR, "Invalid fix reaxff/species nevery value {}", nevery); + if (nrepeat <= 0) error->all(FLERR, "Invalid fix reaxff/species nrepeat value {}", nrepeat); + if (nfreq <= 0) error->all(FLERR, "Invalid fix reaxff/species nfreq value {}", nfreq); + if ((nfreq % nevery) || (nrepeat * nevery > nfreq)) + error->all(FLERR, "Incompatible fix reaxff/species nevery/nrepeat/nfreq settings"); // Neighbor lists must stay unchanged during averaging of bonds, // but may be updated when no averaging is performed. @@ -99,13 +102,13 @@ FixReaxFFSpecies::FixReaxFFSpecies(LAMMPS *lmp, int narg, char **arg) : Fix(lmp, rene_flag = 1; } - if (me == 0 && rene_flag) + if (comm->me == 0 && rene_flag) error->warning(FLERR, "Resetting reneighboring criteria to 'delay {} every {} check no' " - "due to fix reaxff/species", + "due to fix reaxff/species averaging of bond data", neighbor->delay, neighbor->every); - if (me == 0) { + if (comm->me == 0) { if (platform::has_compress_extension(arg[6])) { fp = platform::compressed_write(arg[6]); compressed = 1; @@ -125,17 +128,6 @@ FixReaxFFSpecies::FixReaxFFSpecies(LAMMPS *lmp, int narg, char **arg) : Fix(lmp, memory->create(clusterID, ntmp, "reaxff/species:clusterID"); vector_atom = clusterID; - BOCut = nullptr; - Name = nullptr; - MolName = nullptr; - MolType = nullptr; - NMol = nullptr; - Mol2Spec = nullptr; - nd = nullptr; - molmap = nullptr; - mark = nullptr; - fdel = nullptr; - nmax = 0; setupflag = 0; @@ -160,16 +152,16 @@ FixReaxFFSpecies::FixReaxFFSpecies(LAMMPS *lmp, int narg, char **arg) : Fix(lmp, int iarg = 7; while (iarg < narg) { - // set BO cutoff if (strcmp(arg[iarg], "cutoff") == 0) { - if (iarg + 4 > narg) error->all(FLERR, "Illegal fix reaxff/species command"); + if (iarg + 4 > narg) utils::missing_cmd_args(FLERR, "fix reaxff/species cutoff", error); itype = utils::inumeric(FLERR, arg[iarg + 1], false, lmp); jtype = utils::inumeric(FLERR, arg[iarg + 2], false, lmp); bo_cut = utils::numeric(FLERR, arg[iarg + 3], false, lmp); - if (itype > ntypes || jtype > ntypes) error->all(FLERR, "Illegal fix reaxff/species command"); - if (itype <= 0 || jtype <= 0) error->all(FLERR, "Illegal fix reaxff/species command"); - if (bo_cut > 1.0 || bo_cut < 0.0) error->all(FLERR, "Illegal fix reaxff/species command"); + if ((itype <= 0) || (jtype <= 0) || (itype > ntypes) || (jtype > ntypes)) + error->all(FLERR, "Fix reaxff/species cutoff atom type(s) out of range"); + if ((bo_cut > 1.0) || (bo_cut < 0.0)) + error->all(FLERR, "Fix reaxff/species invalid cutoff value: {}", bo_cut); BOCut[itype][jtype] = bo_cut; BOCut[jtype][itype] = bo_cut; @@ -177,7 +169,8 @@ FixReaxFFSpecies::FixReaxFFSpecies(LAMMPS *lmp, int narg, char **arg) : Fix(lmp, // modify element type names } else if (strcmp(arg[iarg], "element") == 0) { - if (iarg + ntypes + 1 > narg) error->all(FLERR, "Illegal fix reaxff/species command"); + if (iarg + ntypes + 1 > narg) + utils::missing_cmd_args(FLERR, "fix reaxff/species element", error); eletype = (char **) malloc(ntypes * sizeof(char *)); int len; @@ -194,63 +187,66 @@ FixReaxFFSpecies::FixReaxFFSpecies(LAMMPS *lmp, int narg, char **arg) : Fix(lmp, delflag = 1; delete[] filedel; filedel = utils::strdup(arg[iarg + 1]); - if (me == 0) { + if (comm->me == 0) { fdel = fopen(filedel, "w"); if (!fdel) error->one(FLERR, "Cannot open fix reaxff/species delete file {}: {}", filedel, utils::getsyserror()); } - del_opened = 1; if (strcmp(arg[iarg + 2], "masslimit") == 0) { - if (iarg + 5 > narg) error->all(FLERR, "Illegal fix reaxff/species command"); + if (iarg + 5 > narg) utils::missing_cmd_args(FLERR, "fix reaxff/species masslimit", error); masslimitflag = 1; - massmin = atof(arg[iarg + 3]); - massmax = atof(arg[iarg + 4]); + massmin = utils::numeric(FLERR, arg[iarg + 3], false, lmp); + massmax = utils::numeric(FLERR, arg[iarg + 4], false, lmp); iarg += 5; + } else if (strcmp(arg[iarg + 2], "specieslist") == 0) { specieslistflag = 1; - ndelspec = atoi(arg[iarg + 3]); - if (iarg + ndelspec + 4 > narg) error->all(FLERR, "Illegal fix reaxff/species command"); + ndelspec = utils::inumeric(FLERR, arg[iarg + 3], false, lmp); + if (iarg + ndelspec + 4 > narg) + utils::missing_cmd_args(FLERR, "fix reaxff/species delete specieslist", error); del_species.resize(ndelspec); for (int i = 0; i < ndelspec; i++) del_species[i] = arg[iarg + 4 + i]; - if (me == 0) { + if (comm->me == 0) { fprintf(fdel, "Timestep"); for (i = 0; i < ndelspec; i++) fprintf(fdel, "\t%s", del_species[i].c_str()); fprintf(fdel, "\n"); fflush(fdel); } - iarg += ndelspec + 4; + } else - error->all(FLERR, "Illegal fix reaxff/species command"); + error->all(FLERR, "Unknown fix reaxff/species delete option: {}", arg[iarg]); // position of molecules } else if (strcmp(arg[iarg], "position") == 0) { - if (iarg + 3 > narg) error->all(FLERR, "Illegal fix reaxff/species command"); + if (iarg + 3 > narg) utils::missing_cmd_args(FLERR, "fix reaxff/species position", error); posflag = 1; posfreq = utils::inumeric(FLERR, arg[iarg + 1], false, lmp); if (posfreq < nfreq || (posfreq % nfreq != 0)) - error->all(FLERR, "Illegal fix reaxff/species command"); + error->all(FLERR, "Incompatible fix reaxff/species postion frequency {}", posfreq); filepos = new char[255]; strcpy(filepos, arg[iarg + 2]); if (strchr(filepos, '*')) { multipos = 1; } else { - if (me == 0) { + if (comm->me == 0) { pos = fopen(filepos, "w"); - if (pos == nullptr) error->one(FLERR, "Cannot open fix reaxff/species position file"); + if (pos == nullptr) + error->one(FLERR, "Cannot open fix reaxff/species position file: {}", + utils::getsyserror()); } singlepos_opened = 1; multipos = 0; } iarg += 3; } else - error->all(FLERR, "Illegal fix reaxff/species command"); + error->all(FLERR, "Unknown fix reaxff/species keyword: {}", arg[iarg]); } if (!eleflag) { @@ -262,7 +258,7 @@ FixReaxFFSpecies::FixReaxFFSpecies(LAMMPS *lmp, int narg, char **arg) : Fix(lmp, } if (delflag && specieslistflag && masslimitflag) - error->all(FLERR, "Illegal fix reaxff/species command"); + error->all(FLERR, "Incompatible combination fix reaxff/species command options"); vector_nmole = 0; vector_nspec = 0; @@ -287,7 +283,7 @@ FixReaxFFSpecies::~FixReaxFFSpecies() delete[] filepos; delete[] filedel; - if (me == 0) { + if (comm->me == 0) { if (compressed) platform::pclose(fp); else @@ -366,7 +362,7 @@ void FixReaxFFSpecies::init_list(int /*id*/, NeighList *ptr) void FixReaxFFSpecies::post_integrate() { Output_ReaxFF_Bonds(update->ntimestep, fp); - if (me == 0) fflush(fp); + if (comm->me == 0) fflush(fp); } /* ---------------------------------------------------------------------- */ @@ -403,11 +399,11 @@ void FixReaxFFSpecies::Output_ReaxFF_Bonds(bigint ntimestep, FILE * /*fp*/) vector_nmole = Nmole; vector_nspec = Nspec; - if (me == 0 && ntimestep >= 0) WriteFormulas(Nmole, Nspec); + if (comm->me == 0 && ntimestep >= 0) WriteFormulas(Nmole, Nspec); if (posflag && ((ntimestep) % posfreq == 0)) { WritePos(Nmole, Nspec); - if (me == 0) fflush(pos); + if (comm->me == 0) fflush(pos); } if (delflag) DeleteSpecies(Nmole, Nspec); @@ -499,7 +495,7 @@ void FixReaxFFSpecies::FindMolecule() if (!anychange) break; MPI_Allreduce(&loop, &looptot, 1, MPI_INT, MPI_SUM, world); - if (looptot >= 400 * nprocs) break; + if (looptot >= 400 * comm->nprocs) break; } } @@ -530,12 +526,12 @@ void FixReaxFFSpecies::SortMolecule(int &Nmole) return; } if (idlo == ntotal) - if (me == 0) + if (comm->me == 0) error->warning(FLERR, "Atom with cluster ID = maxmol included in fix reaxff/species group {}", group->names[igroup]); MPI_Allreduce(&flag, &flagall, 1, MPI_INT, MPI_SUM, world); - if (flagall && me == 0) + if (flagall && comm->me == 0) error->warning(FLERR, "Atom with cluster ID = 0 included in fix reaxff/species group {}", group->names[igroup]); @@ -712,7 +708,7 @@ void FixReaxFFSpecies::WriteFormulas(int Nmole, int Nspec) void FixReaxFFSpecies::OpenPos() { - if (me == 0) { + if (comm->me == 0) { auto filecurrent = utils::star_subst(filepos, update->ntimestep, padflag); pos = fopen(filecurrent.c_str(), "w"); if (pos == nullptr) @@ -742,7 +738,7 @@ void FixReaxFFSpecies::WritePos(int Nmole, int Nspec) for (int j = 0; j < 3; j++) halfbox[j] = box[j] / 2; - if (me == 0) { + if (comm->me == 0) { fmt::print(pos, "Timestep {} NMole {} NSpec {} xlo {:f} " "xhi {:f} ylo {:f} yhi {:f} zlo {:f} zhi {:f}\n", @@ -796,7 +792,7 @@ void FixReaxFFSpecies::WritePos(int Nmole, int Nspec) MPI_Reduce(Name, Nameall, ntypes, MPI_INT, MPI_SUM, 0, world); for (n = 0; n < ntypes; n++) Name[n] = Nameall[n]; - if (me == 0) { + if (comm->me == 0) { fprintf(pos, "%d\t%d\t", m, count); for (n = 0; n < ntypes; n++) { if (Name[n] != 0) { @@ -822,7 +818,7 @@ void FixReaxFFSpecies::WritePos(int Nmole, int Nspec) fprintf(pos, "\n"); } } - if (me == 0 && !multipos) fprintf(pos, "#\n"); + if (comm->me == 0 && !multipos) fprintf(pos, "#\n"); memory->destroy(Nameall); } @@ -935,12 +931,12 @@ void FixReaxFFSpecies::DeleteSpecies(int Nmole, int Nspec) atom->natoms -= ndel; - if (me == 0) + if (comm->me == 0) MPI_Reduce(MPI_IN_PLACE, deletecount, ndelcomm, MPI_DOUBLE, MPI_SUM, 0, world); else MPI_Reduce(deletecount, deletecount, ndelcomm, MPI_DOUBLE, MPI_SUM, 0, world); - if (me == 0) { + if (comm->me == 0) { if (masslimitflag) { int printflag = 0; for (int m = 0; m < Nspec; m++) { diff --git a/src/REAXFF/fix_reaxff_species.h b/src/REAXFF/fix_reaxff_species.h index e272ff7d6c..78eb39a0b6 100644 --- a/src/REAXFF/fix_reaxff_species.h +++ b/src/REAXFF/fix_reaxff_species.h @@ -44,7 +44,7 @@ class FixReaxFFSpecies : public Fix { double compute_vector(int) override; protected: - int me, nprocs, nmax, nlocal, ntypes, ntotal; + int nmax, nlocal, ntypes, ntotal; int nrepeat, nfreq, posfreq, compressed, ndelspec; int Nmoltype, vector_nmole, vector_nspec; int *Name, *MolName, *NMol, *nd, *MolType, *molmap, *mark; From 0f33ff1fc19ae22745740d934638eadb03e2b797 Mon Sep 17 00:00:00 2001 From: Stan Gerald Moore Date: Mon, 10 Oct 2022 12:08:31 -0600 Subject: [PATCH 4/6] Avoid using a host pointer in device code --- src/KOKKOS/pair_meam_kokkos.cpp | 224 ++++++++++++++++++-------------- src/KOKKOS/pair_meam_kokkos.h | 6 + 2 files changed, 130 insertions(+), 100 deletions(-) diff --git a/src/KOKKOS/pair_meam_kokkos.cpp b/src/KOKKOS/pair_meam_kokkos.cpp index c0f18a2969..86b4a2ed51 100644 --- a/src/KOKKOS/pair_meam_kokkos.cpp +++ b/src/KOKKOS/pair_meam_kokkos.cpp @@ -123,6 +123,7 @@ void PairMEAMKokkos::compute(int eflag_in, int vflag_in) Kokkos::parallel_reduce(Kokkos::RangePolicy(0,inum_half),*this,n); meam_inst_kk->meam_dens_setup(atom->nmax, nall, n); + update_meam_views(); x = atomKK->k_x.view(); f = atomKK->k_f.view(); @@ -324,35 +325,35 @@ KOKKOS_INLINE_FUNCTION void PairMEAMKokkos::operator()(TagPairMEAMPackForwardComm, const int &i) const { int j = d_sendlist(iswap, i); int m = i*38; - v_buf[m++] = meam_inst_kk->d_rho0[j]; - v_buf[m++] = meam_inst_kk->d_rho1[j]; - v_buf[m++] = meam_inst_kk->d_rho2[j]; - v_buf[m++] = meam_inst_kk->d_rho3[j]; - v_buf[m++] = meam_inst_kk->d_frhop[j]; - v_buf[m++] = meam_inst_kk->d_gamma[j]; - v_buf[m++] = meam_inst_kk->d_dgamma1[j]; - v_buf[m++] = meam_inst_kk->d_dgamma2[j]; - v_buf[m++] = meam_inst_kk->d_dgamma3[j]; - v_buf[m++] = meam_inst_kk->d_arho2b[j]; - v_buf[m++] = meam_inst_kk->d_arho1(j,0); - v_buf[m++] = meam_inst_kk->d_arho1(j,1); - v_buf[m++] = meam_inst_kk->d_arho1(j,2); - v_buf[m++] = meam_inst_kk->d_arho2(j,0); - v_buf[m++] = meam_inst_kk->d_arho2(j,1); - v_buf[m++] = meam_inst_kk->d_arho2(j,2); - v_buf[m++] = meam_inst_kk->d_arho2(j,3); - v_buf[m++] = meam_inst_kk->d_arho2(j,4); - v_buf[m++] = meam_inst_kk->d_arho2(j,5); - for (int k = 0; k < 10; k++) v_buf[m++] = meam_inst_kk->d_arho3(j,k); - v_buf[m++] = meam_inst_kk->d_arho3b(j,0); - v_buf[m++] = meam_inst_kk->d_arho3b(j,1); - v_buf[m++] = meam_inst_kk->d_arho3b(j,2); - v_buf[m++] = meam_inst_kk->d_t_ave(j,0); - v_buf[m++] = meam_inst_kk->d_t_ave(j,1); - v_buf[m++] = meam_inst_kk->d_t_ave(j,2); - v_buf[m++] = meam_inst_kk->d_tsq_ave(j,0); - v_buf[m++] = meam_inst_kk->d_tsq_ave(j,1); - v_buf[m++] = meam_inst_kk->d_tsq_ave(j,2); + v_buf[m++] = d_rho0[j]; + v_buf[m++] = d_rho1[j]; + v_buf[m++] = d_rho2[j]; + v_buf[m++] = d_rho3[j]; + v_buf[m++] = d_frhop[j]; + v_buf[m++] = d_gamma[j]; + v_buf[m++] = d_dgamma1[j]; + v_buf[m++] = d_dgamma2[j]; + v_buf[m++] = d_dgamma3[j]; + v_buf[m++] = d_arho2b[j]; + v_buf[m++] = d_arho1(j,0); + v_buf[m++] = d_arho1(j,1); + v_buf[m++] = d_arho1(j,2); + v_buf[m++] = d_arho2(j,0); + v_buf[m++] = d_arho2(j,1); + v_buf[m++] = d_arho2(j,2); + v_buf[m++] = d_arho2(j,3); + v_buf[m++] = d_arho2(j,4); + v_buf[m++] = d_arho2(j,5); + for (int k = 0; k < 10; k++) v_buf[m++] = d_arho3(j,k); + v_buf[m++] = d_arho3b(j,0); + v_buf[m++] = d_arho3b(j,1); + v_buf[m++] = d_arho3b(j,2); + v_buf[m++] = d_t_ave(j,0); + v_buf[m++] = d_t_ave(j,1); + v_buf[m++] = d_t_ave(j,2); + v_buf[m++] = d_tsq_ave(j,0); + v_buf[m++] = d_tsq_ave(j,1); + v_buf[m++] = d_tsq_ave(j,2); } /* ---------------------------------------------------------------------- */ @@ -372,35 +373,35 @@ KOKKOS_INLINE_FUNCTION void PairMEAMKokkos::operator()(TagPairMEAMUnpackForwardComm, const int &i) const{ int m = i*38; - meam_inst_kk->d_rho0[i+first] = v_buf[m++]; - meam_inst_kk->d_rho1[i+first] = v_buf[m++]; - meam_inst_kk->d_rho2[i+first] = v_buf[m++]; - meam_inst_kk->d_rho3[i+first] = v_buf[m++]; - meam_inst_kk->d_frhop[i+first] = v_buf[m++]; - meam_inst_kk->d_gamma[i+first] = v_buf[m++]; - meam_inst_kk->d_dgamma1[i+first] = v_buf[m++]; - meam_inst_kk->d_dgamma2[i+first] = v_buf[m++]; - meam_inst_kk->d_dgamma3[i+first] = v_buf[m++]; - meam_inst_kk->d_arho2b[i+first] = v_buf[m++]; - meam_inst_kk->d_arho1(i+first,0) = v_buf[m++]; - meam_inst_kk->d_arho1(i+first,1) = v_buf[m++]; - meam_inst_kk->d_arho1(i+first,2) = v_buf[m++]; - meam_inst_kk->d_arho2(i+first,0) = v_buf[m++]; - meam_inst_kk->d_arho2(i+first,1) = v_buf[m++]; - meam_inst_kk->d_arho2(i+first,2) = v_buf[m++]; - meam_inst_kk->d_arho2(i+first,3) = v_buf[m++]; - meam_inst_kk->d_arho2(i+first,4) = v_buf[m++]; - meam_inst_kk->d_arho2(i+first,5) = v_buf[m++]; - for (int k = 0; k < 10; k++) meam_inst_kk->d_arho3(i+first,k) = v_buf[m++]; - meam_inst_kk->d_arho3b(i+first,0) = v_buf[m++]; - meam_inst_kk->d_arho3b(i+first,1) = v_buf[m++]; - meam_inst_kk->d_arho3b(i+first,2) = v_buf[m++]; - meam_inst_kk->d_t_ave(i+first,0) = v_buf[m++]; - meam_inst_kk->d_t_ave(i+first,1) = v_buf[m++]; - meam_inst_kk->d_t_ave(i+first,2) = v_buf[m++]; - meam_inst_kk->d_tsq_ave(i+first,0) = v_buf[m++]; - meam_inst_kk->d_tsq_ave(i+first,1) = v_buf[m++]; - meam_inst_kk->d_tsq_ave(i+first,2) = v_buf[m++]; + d_rho0[i+first] = v_buf[m++]; + d_rho1[i+first] = v_buf[m++]; + d_rho2[i+first] = v_buf[m++]; + d_rho3[i+first] = v_buf[m++]; + d_frhop[i+first] = v_buf[m++]; + d_gamma[i+first] = v_buf[m++]; + d_dgamma1[i+first] = v_buf[m++]; + d_dgamma2[i+first] = v_buf[m++]; + d_dgamma3[i+first] = v_buf[m++]; + d_arho2b[i+first] = v_buf[m++]; + d_arho1(i+first,0) = v_buf[m++]; + d_arho1(i+first,1) = v_buf[m++]; + d_arho1(i+first,2) = v_buf[m++]; + d_arho2(i+first,0) = v_buf[m++]; + d_arho2(i+first,1) = v_buf[m++]; + d_arho2(i+first,2) = v_buf[m++]; + d_arho2(i+first,3) = v_buf[m++]; + d_arho2(i+first,4) = v_buf[m++]; + d_arho2(i+first,5) = v_buf[m++]; + for (int k = 0; k < 10; k++) d_arho3(i+first,k) = v_buf[m++]; + d_arho3b(i+first,0) = v_buf[m++]; + d_arho3b(i+first,1) = v_buf[m++]; + d_arho3b(i+first,2) = v_buf[m++]; + d_t_ave(i+first,0) = v_buf[m++]; + d_t_ave(i+first,1) = v_buf[m++]; + d_t_ave(i+first,2) = v_buf[m++]; + d_tsq_ave(i+first,0) = v_buf[m++]; + d_tsq_ave(i+first,1) = v_buf[m++]; + d_tsq_ave(i+first,2) = v_buf[m++]; } /* ---------------------------------------------------------------------- */ @@ -555,27 +556,27 @@ KOKKOS_INLINE_FUNCTION void PairMEAMKokkos::operator()(TagPairMEAMPackReverseComm, const int &i) const { int m = i*30; - v_buf[m++] = meam_inst_kk->d_rho0[i+first]; - v_buf[m++] = meam_inst_kk->d_arho2b[i+first]; - v_buf[m++] = meam_inst_kk->d_arho1(i+first,0); - v_buf[m++] = meam_inst_kk->d_arho1(i+first,1); - v_buf[m++] = meam_inst_kk->d_arho1(i+first,2); - v_buf[m++] = meam_inst_kk->d_arho2(i+first,0); - v_buf[m++] = meam_inst_kk->d_arho2(i+first,1); - v_buf[m++] = meam_inst_kk->d_arho2(i+first,2); - v_buf[m++] = meam_inst_kk->d_arho2(i+first,3); - v_buf[m++] = meam_inst_kk->d_arho2(i+first,4); - v_buf[m++] = meam_inst_kk->d_arho2(i+first,5); - for (int k = 0; k < 10; k++) v_buf[m++] = meam_inst_kk->d_arho3(i+first,k); - v_buf[m++] = meam_inst_kk->d_arho3b(i+first,0); - v_buf[m++] = meam_inst_kk->d_arho3b(i+first,1); - v_buf[m++] = meam_inst_kk->d_arho3b(i+first,2); - v_buf[m++] = meam_inst_kk->d_t_ave(i+first,0); - v_buf[m++] = meam_inst_kk->d_t_ave(i+first,1); - v_buf[m++] = meam_inst_kk->d_t_ave(i+first,2); - v_buf[m++] = meam_inst_kk->d_tsq_ave(i+first,0); - v_buf[m++] = meam_inst_kk->d_tsq_ave(i+first,1); - v_buf[m++] = meam_inst_kk->d_tsq_ave(i+first,2); + v_buf[m++] = d_rho0[i+first]; + v_buf[m++] = d_arho2b[i+first]; + v_buf[m++] = d_arho1(i+first,0); + v_buf[m++] = d_arho1(i+first,1); + v_buf[m++] = d_arho1(i+first,2); + v_buf[m++] = d_arho2(i+first,0); + v_buf[m++] = d_arho2(i+first,1); + v_buf[m++] = d_arho2(i+first,2); + v_buf[m++] = d_arho2(i+first,3); + v_buf[m++] = d_arho2(i+first,4); + v_buf[m++] = d_arho2(i+first,5); + for (int k = 0; k < 10; k++) v_buf[m++] = d_arho3(i+first,k); + v_buf[m++] = d_arho3b(i+first,0); + v_buf[m++] = d_arho3b(i+first,1); + v_buf[m++] = d_arho3b(i+first,2); + v_buf[m++] = d_t_ave(i+first,0); + v_buf[m++] = d_t_ave(i+first,1); + v_buf[m++] = d_t_ave(i+first,2); + v_buf[m++] = d_tsq_ave(i+first,0); + v_buf[m++] = d_tsq_ave(i+first,1); + v_buf[m++] = d_tsq_ave(i+first,2); } /* ---------------------------------------------------------------------- */ @@ -640,27 +641,27 @@ void PairMEAMKokkos::operator()(TagPairMEAMUnpackReverseComm, const int j = d_sendlist(iswap, i); int m = i*30; - meam_inst_kk->d_rho0[j] += v_buf[m++]; - meam_inst_kk->d_arho2b[j] += v_buf[m++]; - meam_inst_kk->d_arho1(j,0) += v_buf[m++]; - meam_inst_kk->d_arho1(j,1) += v_buf[m++]; - meam_inst_kk->d_arho1(j,2) += v_buf[m++]; - meam_inst_kk->d_arho2(j,0) += v_buf[m++]; - meam_inst_kk->d_arho2(j,1) += v_buf[m++]; - meam_inst_kk->d_arho2(j,2) += v_buf[m++]; - meam_inst_kk->d_arho2(j,3) += v_buf[m++]; - meam_inst_kk->d_arho2(j,4) += v_buf[m++]; - meam_inst_kk->d_arho2(j,5) += v_buf[m++]; - for (int k = 0; k < 10; k++) meam_inst_kk->d_arho3(j,k) += v_buf[m++]; - meam_inst_kk->d_arho3b(j,0) += v_buf[m++]; - meam_inst_kk->d_arho3b(j,1) += v_buf[m++]; - meam_inst_kk->d_arho3b(j,2) += v_buf[m++]; - meam_inst_kk->d_t_ave(j,0) += v_buf[m++]; - meam_inst_kk->d_t_ave(j,1) += v_buf[m++]; - meam_inst_kk->d_t_ave(j,2) += v_buf[m++]; - meam_inst_kk->d_tsq_ave(j,0) += v_buf[m++]; - meam_inst_kk->d_tsq_ave(j,1) += v_buf[m++]; - meam_inst_kk->d_tsq_ave(j,2) += v_buf[m++]; + d_rho0[j] += v_buf[m++]; + d_arho2b[j] += v_buf[m++]; + d_arho1(j,0) += v_buf[m++]; + d_arho1(j,1) += v_buf[m++]; + d_arho1(j,2) += v_buf[m++]; + d_arho2(j,0) += v_buf[m++]; + d_arho2(j,1) += v_buf[m++]; + d_arho2(j,2) += v_buf[m++]; + d_arho2(j,3) += v_buf[m++]; + d_arho2(j,4) += v_buf[m++]; + d_arho2(j,5) += v_buf[m++]; + for (int k = 0; k < 10; k++) d_arho3(j,k) += v_buf[m++]; + d_arho3b(j,0) += v_buf[m++]; + d_arho3b(j,1) += v_buf[m++]; + d_arho3b(j,2) += v_buf[m++]; + d_t_ave(j,0) += v_buf[m++]; + d_t_ave(j,1) += v_buf[m++]; + d_t_ave(j,2) += v_buf[m++]; + d_tsq_ave(j,0) += v_buf[m++]; + d_tsq_ave(j,1) += v_buf[m++]; + d_tsq_ave(j,2) += v_buf[m++]; } /* ---------------------------------------------------------------------- */ @@ -744,6 +745,29 @@ void PairMEAMKokkos::operator()(TagPairMEAMOffsets, const int ii, in /* ---------------------------------------------------------------------- */ +template +void PairMEAMKokkos::update_meam_views() +{ + d_rho0 = meam_inst_kk->d_rho0; + d_rho1 = meam_inst_kk->d_rho1; + d_rho2 = meam_inst_kk->d_rho2; + d_rho3 = meam_inst_kk->d_rho3; + d_frhop = meam_inst_kk->d_frhop; + d_gamma = meam_inst_kk->d_gamma; + d_dgamma1 = meam_inst_kk->d_dgamma1; + d_dgamma2 = meam_inst_kk->d_dgamma2; + d_dgamma3 = meam_inst_kk->d_dgamma3; + d_arho1 = meam_inst_kk->d_arho1; + d_arho2 = meam_inst_kk->d_arho2; + d_arho3 = meam_inst_kk->d_arho3; + d_arho2b = meam_inst_kk->d_arho2b; + d_arho3b = meam_inst_kk->d_arho3b; + d_t_ave = meam_inst_kk->d_t_ave; + d_tsq_ave = meam_inst_kk->d_tsq_ave; +} + +/* ---------------------------------------------------------------------- */ + namespace LAMMPS_NS { template class PairMEAMKokkos; #ifdef KOKKOS_ENABLE_CUDA diff --git a/src/KOKKOS/pair_meam_kokkos.h b/src/KOKKOS/pair_meam_kokkos.h index fef4266b8a..a213c1e9b1 100644 --- a/src/KOKKOS/pair_meam_kokkos.h +++ b/src/KOKKOS/pair_meam_kokkos.h @@ -114,6 +114,12 @@ class PairMEAMKokkos : public PairMEAM, public KokkosBase { int iswap,first; int neighflag,nlocal,nall,eflag,vflag; + typename ArrayTypes::t_ffloat_1d d_rho, d_rho0, d_rho1, d_rho2, d_rho3, d_frhop; + typename ArrayTypes::t_ffloat_1d d_gamma, d_dgamma1, d_dgamma2, d_dgamma3, d_arho2b; + typename ArrayTypes::t_ffloat_2d d_arho1, d_arho2, d_arho3, d_arho3b, d_t_ave, d_tsq_ave; + + void update_meam_views(); + friend void pair_virial_fdotr_compute(PairMEAMKokkos*); }; From 185639f077404e10b840f5e32e8759e231c0d91b Mon Sep 17 00:00:00 2001 From: Axel Kohlmeyer Date: Mon, 10 Oct 2022 16:37:04 -0400 Subject: [PATCH 5/6] must use utils::inumeric() for parsing integers --- src/input.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/src/input.cpp b/src/input.cpp index 09e5d6decb..ee79bd0d15 100644 --- a/src/input.cpp +++ b/src/input.cpp @@ -1740,9 +1740,9 @@ void Input::pair_coeff() // if arg[1] < arg[0], and neither contain a wildcard, reorder int itype,jtype; - if (strchr(arg[0],'*') == nullptr && strchr(arg[1],'*') == nullptr) { - itype = utils::numeric(FLERR,arg[0],false,lmp); - jtype = utils::numeric(FLERR,arg[1],false,lmp); + if (utils::strmatch(arg[0],"^\\d+$") && utils::strmatch(arg[1],"^\\d+$")) { + itype = utils::inumeric(FLERR,arg[0],false,lmp); + jtype = utils::inumeric(FLERR,arg[1],false,lmp); if (jtype < itype) { char *str = arg[0]; arg[0] = arg[1]; From 5c82c8a8a3ebec6076fc95b3ecaec3275eb52a99 Mon Sep 17 00:00:00 2001 From: Axel Kohlmeyer Date: Mon, 10 Oct 2022 16:37:20 -0400 Subject: [PATCH 6/6] improve error messages --- src/input.cpp | 9 ++++----- 1 file changed, 4 insertions(+), 5 deletions(-) diff --git a/src/input.cpp b/src/input.cpp index ee79bd0d15..0298f14ef1 100644 --- a/src/input.cpp +++ b/src/input.cpp @@ -1726,11 +1726,10 @@ void Input::pair_coeff() { if (domain->box_exist == 0) error->all(FLERR,"Pair_coeff command before simulation box is defined"); - if (force->pair == nullptr) - error->all(FLERR,"Pair_coeff command before pair_style is defined"); - if ((narg < 2) || (force->pair->one_coeff && ((strcmp(arg[0],"*") != 0) - || (strcmp(arg[1],"*") != 0)))) - error->all(FLERR,"Incorrect args for pair coefficients"); + if (force->pair == nullptr) error->all(FLERR,"Pair_coeff command without a pair style"); + if (narg < 2) utils::missing_cmd_args(FLERR,"pair_coeff", error); + if (force->pair->one_coeff && ((strcmp(arg[0],"*") != 0) || (strcmp(arg[1],"*") != 0))) + error->all(FLERR,"Pair_coeff must start with * * for this pair style"); char *newarg0 = utils::expand_type(FLERR, arg[0], Atom::ATOM, lmp); if (newarg0) arg[0] = newarg0;