Page Menu
Home
c4science
Search
Configure Global Search
Log In
Files
F110810965
neighbor_cuda.cpp
No One
Temporary
Actions
Download File
Edit File
Delete File
View Transforms
Subscribe
Mute Notifications
Award Token
Subscribers
None
File Metadata
Details
File Info
Storage
Attached
Created
Mon, Apr 28, 05:42
Size
6 KB
Mime Type
text/x-c
Expires
Wed, Apr 30, 05:42 (2 d)
Engine
blob
Format
Raw Data
Handle
25846122
Attached To
rLAMMPS lammps
neighbor_cuda.cpp
View Options
/* ----------------------------------------------------------------------
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.
------------------------------------------------------------------------- */
#include "math.h"
#include "neighbor_cuda.h"
#include "cuda.h"
#include "atom.h"
#include "atom_vec.h"
#include "domain.h"
#include "neigh_list.h"
#include "neigh_request.h"
#include "force.h"
#include "group.h"
#include "memory.h"
#include "error.h"
using namespace LAMMPS_NS;
#define MIN(a,b) ((a) < (b) ? (a) : (b))
#define MAX(a,b) ((a) > (b) ? (a) : (b))
enum{NSQ,BIN,MULTI}; // also in neigh_list.cpp
/* ---------------------------------------------------------------------- */
NeighborCuda::NeighborCuda(LAMMPS *lmp) : Neighbor(lmp)
{
cuda = lmp->cuda;
if(cuda == NULL)
error->all("You cannot use a /cuda class, without activating 'cuda' acceleration. Provide '-c on' as command-line argument to LAMMPS..");
}
/* ---------------------------------------------------------------------- */
void NeighborCuda::init()
{
cuda->set_neighinit(dist_check,0.25*skin*skin);
cudable = 1;
Neighbor::init();
}
/* ----------------------------------------------------------------------
overwrite either full_nsq or full_bin with CUDA-equivalent methods
any other neighbor build method is unchanged
------------------------------------------------------------------------- */
void NeighborCuda::choose_build(int index, NeighRequest *rq)
{
Neighbor::choose_build(index,rq);
if (rq->full && style == NSQ && rq->ghost == 0 && rq->cudable)
pair_build[index] = (Neighbor::PairPtr) &NeighborCuda::full_nsq_cuda;
else if (rq->full && style == BIN && rq->ghost == 0 && rq->cudable)
pair_build[index] = (Neighbor::PairPtr) &NeighborCuda::full_bin_cuda;
}
/* ---------------------------------------------------------------------- */
int NeighborCuda::check_distance()
{
double delx,dely,delz,rsq;
double delta,deltasq,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;
double **x = atom->x;
int nlocal = atom->nlocal;
if (includegroup) nlocal = atom->nfirst;
int flag = 0;
if (not cuda->neighbor_decide_by_integrator) {
cuda->cu_x_download();
for (int i = 0; i < nlocal; i++) {
delx = x[i][0] - xhold[i][0];
dely = x[i][1] - xhold[i][1];
delz = x[i][2] - xhold[i][2];
rsq = delx*delx + dely*dely + delz*delz;
if (rsq > deltasq) flag = 1;
}
}
else flag = cuda->shared_data.atom.reneigh_flag;
int flagall;
MPI_Allreduce(&flag,&flagall,1,MPI_INT,MPI_MAX,world);
if (flagall && ago == MAX(every,delay)) ndanger++;
return flagall;
}
/* ---------------------------------------------------------------------- */
void NeighborCuda::build()
{
int i;
ago = 0;
ncalls++;
// store current atom positions and box size if needed
if (dist_check) {
if (cuda->decide_by_integrator())
cuda->update_xhold(maxhold, &xhold[0][0]);
else {
if (cuda->finished_setup) cuda->cu_x_download();
double **x = atom->x;
int nlocal = atom->nlocal;
if (includegroup) nlocal = atom->nfirst;
if (nlocal > maxhold) {
maxhold = atom->nmax;
memory->destroy(xhold);
memory->create(xhold,maxhold,3,"neigh:xhold");
}
for (i = 0; i < nlocal; i++) {
xhold[i][0] = x[i][0];
xhold[i][1] = x[i][1];
xhold[i][2] = x[i][2];
}
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 (not cudable && cuda->finished_setup && atom->avec->cudable)
cuda->downloadAll();
if (cudable && (not cuda->finished_setup)) {
cuda->checkResize();
cuda->uploadAll();
}
// 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 done for lists with growflag set and which are perpetual
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 neighbor list with special bond flags will not overflow
if (atom->nlocal+atom->nghost > NEIGHMASK)
error->one("Too many local+ghost atoms for neighbor list");
// invoke building of pair and molecular neighbor lists
// only for pairwise lists with buildflag set
for (i = 0; i < nblist; i++)
(this->*pair_build[blist[i]])(lists[blist[i]]);
if (atom->molecular) {
if (force->bond) (this->*bond_build)();
if (force->angle) (this->*angle_build)();
if (force->dihedral) (this->*dihedral_build)();
if (force->improper) (this->*improper_build)();
}
}
Event Timeline
Log In to Comment