Homec4science
Diffusion lammps (master) Browse

lammps/src/KOKKOS/pair_lj_cut_kokkos_light.cpp
4b755fb93b07master

Last Change
Raw File

pair_lj_cut_kokkos_light.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.
------------------------------------------------------------------------- */
/* ----------------------------------------------------------------------
Contributing author: Paul Crozier (SNL)
------------------------------------------------------------------------- */
#include <math.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include "pair_lj_cut_kokkos_light.h"
#include "atom.h"
#include "comm.h"
#include "force.h"
#include "neighbor.h"
#include "neigh_list.h"
#include "neigh_request.h"
#include "update.h"
#include "integrate.h"
#include "respa.h"
#include "math_const.h"
#include "memory.h"
#include "error.h"
#include "kokkos.h"
#include "atom_kokkos.h"
#include "atom_masks.h"
using namespace LAMMPS_NS;
using namespace MathConst;
/* ---------------------------------------------------------------------- */
PairLJCutKokkosLight::PairLJCutKokkosLight(LAMMPS *lmp) : Pair(lmp)
{
respa_enable = 1;
writedata = 1;
atomKK = (AtomKokkos *) atom;
execution_space = Device;
datamask_read = X_MASK | F_MASK | TYPE_MASK | ENERGY_MASK | VIRIAL_MASK;
datamask_modify = F_MASK | ENERGY_MASK | VIRIAL_MASK;
cutsq = NULL;
}
/* ---------------------------------------------------------------------- */
PairLJCutKokkosLight::~PairLJCutKokkosLight()
{
if (allocated) {
memory->destroy(setflag);
memory->destroy(cutsq);
}
}
/* ---------------------------------------------------------------------- */
void PairLJCutKokkosLight::compute(int eflag, int vflag)
{
if (eflag || vflag) ev_setup(eflag,vflag);
else evflag = vflag_fdotr = 0;
Kokkos::View<const double*[3]> x = atomKK->k_x.d_view;;
Kokkos::View<double*[3],Kokkos::MemoryTraits<Kokkos::Atomic> > f = atomKK->k_f.d_view;;
Kokkos::View<const int*> type = atomKK->k_type.d_view;
int nlocal = atom->nlocal;
double special_lj[4];
special_lj[0] = force->special_lj[0]; special_lj[1] = force->special_lj[1];
special_lj[2] = force->special_lj[2]; special_lj[3] = force->special_lj[3];
int newton_pair = force->newton_pair;
NeighListKokkos<Kokkos::DefaultExecutionSpace> kk_list = *((NeighListKokkos<Kokkos::DefaultExecutionSpace>* )list);
const int inum = kk_list.inum;
Kokkos::View<const int*> ilist = kk_list.d_ilist;
Kokkos::View<const int*> numneigh = kk_list.d_numneigh;
copymode = 1;
// loop over neighbors of my atoms
Kokkos::parallel_for( inum, KOKKOS_LAMBDA(const int& ii) {
const int i = ilist(ii);
const double xtmp = x(i,0);
const double ytmp = x(i,0);
const double ztmp = x(i,0);
const int itype = type(i);
const AtomNeighborsConst jlist = kk_list.get_neighbors_const(i);
const int jnum = numneigh(i);
double fx = 0.0;
double fy = 0.0;
double fz = 0.0;
for (int jj = 0; jj < jnum; jj++) {
int j = jlist(jj);
const double factor_lj = special_lj[sbmask(j)];
j &= NEIGHMASK;
const double delx = xtmp - x(j,0);
const double dely = ytmp - x(j,1);
const double delz = ztmp - x(j,2);
const double rsq = delx*delx + dely*dely + delz*delz;
const int jtype = type(j);
if (rsq < cut2(itype,jtype)) {
const double r2inv = 1.0/rsq;
const double r6inv = r2inv*r2inv*r2inv;
const double forcelj = r6inv * (lj1(itype,jtype)*r6inv - lj2(itype,jtype));
const double fpair = factor_lj*forcelj*r2inv;
fx += delx*fpair;
fy += dely*fpair;
fz += delz*fpair;
if (newton_pair || j < nlocal) {
f(j,0) -= delx*fpair;
f(j,1) -= dely*fpair;
f(j,2) -= delz*fpair;
}
if (eflag) {
double evdwl = r6inv*(lj3(itype,jtype)*r6inv-lj4(itype,jtype)) -
offset(itype,jtype);
evdwl *= factor_lj;
if (vflag_either || eflag_atom) ev_tally(ev,i,j,evdwl,fpair,delx,dely,delz);
}
}
}
f(i,0) += fx;
f(i,1) += fy;
f(i,2) += fz;
});
copymode = 0;
if (vflag_fdotr) pair_virial_fdotr_compute(this);
}
/* ---------------------------------------------------------------------- */
void PairLJCutKokkosLight::compute_inner()
{
/*int i,j,ii,jj,inum,jnum,itype,jtype;
double xtmp,ytmp,ztmp,delx,dely,delz,fpair;
double rsq,r2inv,r6inv,forcelj,factor_lj,rsw;
int *ilist,*jlist,*numneigh,**firstneigh;
double **x = atom->x;
double **f = atom->f;
int *type = atom->type;
int nlocal = atom->nlocal;
double *special_lj = force->special_lj;
int newton_pair = force->newton_pair;
inum = listinner->inum;
ilist = listinner->ilist;
numneigh = listinner->numneigh;
firstneigh = listinner->firstneigh;
double cut_out_on = cut_respa[0];
double cut_out_off = cut_respa[1];
double cut_out_diff = cut_out_off - cut_out_on;
double cut_out_on_sq = cut_out_on*cut_out_on;
double cut_out_off_sq = cut_out_off*cut_out_off;
// loop over neighbors of my atoms
for (ii = 0; ii < inum; ii++) {
i = ilist[ii];
xtmp = x[i][0];
ytmp = x[i][1];
ztmp = x[i][2];
itype = type[i];
jlist = firstneigh[i];
jnum = numneigh[i];
for (jj = 0; jj < jnum; jj++) {
j = jlist[jj];
factor_lj = special_lj[sbmask(j)];
j &= NEIGHMASK;
delx = xtmp - x[j][0];
dely = ytmp - x[j][1];
delz = ztmp - x[j][2];
rsq = delx*delx + dely*dely + delz*delz;
if (rsq < cut_out_off_sq) {
r2inv = 1.0/rsq;
r6inv = r2inv*r2inv*r2inv;
jtype = type[j];
forcelj = r6inv * (lj1[itype][jtype]*r6inv - lj2[itype][jtype]);
fpair = factor_lj*forcelj*r2inv;
if (rsq > cut_out_on_sq) {
rsw = (sqrt(rsq) - cut_out_on)/cut_out_diff;
fpair *= 1.0 - rsw*rsw*(3.0 - 2.0*rsw);
}
f[i][0] += delx*fpair;
f[i][1] += dely*fpair;
f[i][2] += delz*fpair;
if (newton_pair || j < nlocal) {
f[j][0] -= delx*fpair;
f[j][1] -= dely*fpair;
f[j][2] -= delz*fpair;
}
}
}
}*/
}
/* ---------------------------------------------------------------------- */
void PairLJCutKokkosLight::compute_middle()
{
/*int i,j,ii,jj,inum,jnum,itype,jtype;
double xtmp,ytmp,ztmp,delx,dely,delz,fpair;
double rsq,r2inv,r6inv,forcelj,factor_lj,rsw;
int *ilist,*jlist,*numneigh,**firstneigh;
double **x = atom->x;
double **f = atom->f;
int *type = atom->type;
int nlocal = atom->nlocal;
double *special_lj = force->special_lj;
int newton_pair = force->newton_pair;
inum = listmiddle->inum;
ilist = listmiddle->ilist;
numneigh = listmiddle->numneigh;
firstneigh = listmiddle->firstneigh;
double cut_in_off = cut_respa[0];
double cut_in_on = cut_respa[1];
double cut_out_on = cut_respa[2];
double cut_out_off = cut_respa[3];
double cut_in_diff = cut_in_on - cut_in_off;
double cut_out_diff = cut_out_off - cut_out_on;
double cut_in_off_sq = cut_in_off*cut_in_off;
double cut_in_on_sq = cut_in_on*cut_in_on;
double cut_out_on_sq = cut_out_on*cut_out_on;
double cut_out_off_sq = cut_out_off*cut_out_off;
// loop over neighbors of my atoms
for (ii = 0; ii < inum; ii++) {
i = ilist[ii];
xtmp = x[i][0];
ytmp = x[i][1];
ztmp = x[i][2];
itype = type[i];
jlist = firstneigh[i];
jnum = numneigh[i];
for (jj = 0; jj < jnum; jj++) {
j = jlist[jj];
factor_lj = special_lj[sbmask(j)];
j &= NEIGHMASK;
delx = xtmp - x[j][0];
dely = ytmp - x[j][1];
delz = ztmp - x[j][2];
rsq = delx*delx + dely*dely + delz*delz;
if (rsq < cut_out_off_sq && rsq > cut_in_off_sq) {
r2inv = 1.0/rsq;
r6inv = r2inv*r2inv*r2inv;
jtype = type[j];
forcelj = r6inv * (lj1[itype][jtype]*r6inv - lj2[itype][jtype]);
fpair = factor_lj*forcelj*r2inv;
if (rsq < cut_in_on_sq) {
rsw = (sqrt(rsq) - cut_in_off)/cut_in_diff;
fpair *= rsw*rsw*(3.0 - 2.0*rsw);
}
if (rsq > cut_out_on_sq) {
rsw = (sqrt(rsq) - cut_out_on)/cut_out_diff;
fpair *= 1.0 + rsw*rsw*(2.0*rsw - 3.0);
}
f[i][0] += delx*fpair;
f[i][1] += dely*fpair;
f[i][2] += delz*fpair;
if (newton_pair || j < nlocal) {
f[j][0] -= delx*fpair;
f[j][1] -= dely*fpair;
f[j][2] -= delz*fpair;
}
}
}
}*/
}
/* ---------------------------------------------------------------------- */
void PairLJCutKokkosLight::compute_outer(int eflag, int vflag)
{
/* int i,j,ii,jj,inum,jnum,itype,jtype;
double xtmp,ytmp,ztmp,delx,dely,delz,evdwl,fpair;
double rsq,r2inv,r6inv,forcelj,factor_lj,rsw;
int *ilist,*jlist,*numneigh,**firstneigh;
evdwl = 0.0;
if (eflag || vflag) ev_setup(eflag,vflag);
else evflag = 0;
double **x = atom->x;
double **f = atom->f;
int *type = atom->type;
int nlocal = atom->nlocal;
double *special_lj = force->special_lj;
int newton_pair = force->newton_pair;
inum = listouter->inum;
ilist = listouter->ilist;
numneigh = listouter->numneigh;
firstneigh = listouter->firstneigh;
double cut_in_off = cut_respa[2];
double cut_in_on = cut_respa[3];
double cut_in_diff = cut_in_on - cut_in_off;
double cut_in_off_sq = cut_in_off*cut_in_off;
double cut_in_on_sq = cut_in_on*cut_in_on;
// loop over neighbors of my atoms
for (ii = 0; ii < inum; ii++) {
i = ilist[ii];
xtmp = x[i][0];
ytmp = x[i][1];
ztmp = x[i][2];
itype = type[i];
jlist = firstneigh[i];
jnum = numneigh[i];
for (jj = 0; jj < jnum; jj++) {
j = jlist[jj];
factor_lj = special_lj[sbmask(j)];
j &= NEIGHMASK;
delx = xtmp - x[j][0];
dely = ytmp - x[j][1];
delz = ztmp - x[j][2];
rsq = delx*delx + dely*dely + delz*delz;
jtype = type[j];
if (rsq < cutsq[itype][jtype]) {
if (rsq > cut_in_off_sq) {
r2inv = 1.0/rsq;
r6inv = r2inv*r2inv*r2inv;
forcelj = r6inv * (lj1[itype][jtype]*r6inv - lj2[itype][jtype]);
fpair = factor_lj*forcelj*r2inv;
if (rsq < cut_in_on_sq) {
rsw = (sqrt(rsq) - cut_in_off)/cut_in_diff;
fpair *= rsw*rsw*(3.0 - 2.0*rsw);
}
f[i][0] += delx*fpair;
f[i][1] += dely*fpair;
f[i][2] += delz*fpair;
if (newton_pair || j < nlocal) {
f[j][0] -= delx*fpair;
f[j][1] -= dely*fpair;
f[j][2] -= delz*fpair;
}
}
if (eflag) {
r2inv = 1.0/rsq;
r6inv = r2inv*r2inv*r2inv;
evdwl = r6inv*(lj3[itype][jtype]*r6inv-lj4[itype][jtype]) -
offset[itype][jtype];
evdwl *= factor_lj;
}
if (vflag) {
if (rsq <= cut_in_off_sq) {
r2inv = 1.0/rsq;
r6inv = r2inv*r2inv*r2inv;
forcelj = r6inv * (lj1[itype][jtype]*r6inv - lj2[itype][jtype]);
fpair = factor_lj*forcelj*r2inv;
} else if (rsq < cut_in_on_sq)
fpair = factor_lj*forcelj*r2inv;
}
if (evflag) ev_tally(i,j,nlocal,newton_pair,
evdwl,0.0,fpair,delx,dely,delz);
}
}
}*/
}
/* ----------------------------------------------------------------------
allocate all arrays
------------------------------------------------------------------------- */
void PairLJCutKokkosLight::allocate()
{
allocated = 1;
int n = atom->ntypes;
memory->create(setflag,n+1,n+1,"pair:setflag");
for (int i = 1; i <= n; i++)
for (int j = i; j <= n; j++)
setflag[i][j] = 0;
memory->create(cutsq,n+1,n+1,"pair:cutsq");
cut = Kokkos::View<double**>("pair:cut",n+1,n+1);
cut2 = Kokkos::View<double**>("pair:cut2",n+1,n+1);
epsilon = Kokkos::View<double**>("pair:epsilon",n+1,n+1);
sigma = Kokkos::View<double**>("pair:sigma",n+1,n+1);
lj1 = Kokkos::View<double**>("pair:lj1",n+1,n+1);
lj2 = Kokkos::View<double**>("pair:lj2",n+1,n+1);
lj3 = Kokkos::View<double**>("pair:lj3",n+1,n+1);
lj4 = Kokkos::View<double**>("pair:lj4",n+1,n+1);
offset = Kokkos::View<double**>("pair:offset",n+1,n+1);
}
/* ----------------------------------------------------------------------
global settings
------------------------------------------------------------------------- */
void PairLJCutKokkosLight::settings(int narg, char **arg)
{
if (narg != 1) error->all(FLERR,"Illegal pair_style command");
cut_global = force->numeric(FLERR,arg[0]);
// reset cutoffs that have been explicitly set
if (allocated) {
int i,j;
for (i = 1; i <= atom->ntypes; i++)
for (j = i+1; j <= atom->ntypes; j++)
if (setflag[i][j]) cut(i,j) = cut_global;
}
}
/* ----------------------------------------------------------------------
set coeffs for one or more type pairs
------------------------------------------------------------------------- */
void PairLJCutKokkosLight::coeff(int narg, char **arg)
{
if (narg < 4 || narg > 5)
error->all(FLERR,"Incorrect args for pair coefficients");
if (!allocated) allocate();
int ilo,ihi,jlo,jhi;
force->bounds(arg[0],atom->ntypes,ilo,ihi);
force->bounds(arg[1],atom->ntypes,jlo,jhi);
double epsilon_one = force->numeric(FLERR,arg[2]);
double sigma_one = force->numeric(FLERR,arg[3]);
double cut_one = cut_global;
if (narg == 5) cut_one = force->numeric(FLERR,arg[4]);
int count = 0;
for (int i = ilo; i <= ihi; i++) {
for (int j = MAX(jlo,i); j <= jhi; j++) {
epsilon(i,j) = epsilon_one;
sigma(i,j) = sigma_one;
cut(i,j) = cut_one;
cut2(i,j) = cut_one*cut_one;
cut2(j,i) = cut_one*cut_one;
setflag[i][j] = 1;
count++;
}
}
if (count == 0) error->all(FLERR,"Incorrect args for pair coefficients");
}
/* ----------------------------------------------------------------------
init specific to this pair style
------------------------------------------------------------------------- */
void PairLJCutKokkosLight::init_style()
{
// request regular or rRESPA neighbor lists
int irequest;
if (update->whichflag == 1 && strstr(update->integrate_style,"respa")) {
int respa = 0;
if (((Respa *) update->integrate)->level_inner >= 0) respa = 1;
if (((Respa *) update->integrate)->level_middle >= 0) respa = 2;
if (respa == 0) irequest = neighbor->request(this,instance_me);
else if (respa == 1) {
irequest = neighbor->request(this,instance_me);
neighbor->requests[irequest]->id = 1;
neighbor->requests[irequest]->half = 0;
neighbor->requests[irequest]->respainner = 1;
irequest = neighbor->request(this,instance_me);
neighbor->requests[irequest]->id = 3;
neighbor->requests[irequest]->half = 0;
neighbor->requests[irequest]->respaouter = 1;
} else {
irequest = neighbor->request(this,instance_me);
neighbor->requests[irequest]->id = 1;
neighbor->requests[irequest]->half = 0;
neighbor->requests[irequest]->respainner = 1;
irequest = neighbor->request(this,instance_me);
neighbor->requests[irequest]->id = 2;
neighbor->requests[irequest]->half = 0;
neighbor->requests[irequest]->respamiddle = 1;
irequest = neighbor->request(this,instance_me);
neighbor->requests[irequest]->id = 3;
neighbor->requests[irequest]->half = 0;
neighbor->requests[irequest]->respaouter = 1;
}
} else irequest = neighbor->request(this,instance_me);
// set rRESPA cutoffs
/*if (strstr(update->integrate_style,"respa") &&
((Respa *) update->integrate)->level_inner >= 0)
cut_respa = ((Respa *) update->integrate)->cutoff;
else*/ cut_respa = NULL;
neighbor->requests[irequest]->
kokkos_host = 0;
neighbor->requests[irequest]->
kokkos_device = 1;
neighbor->requests[irequest]->full = 0;
neighbor->requests[irequest]->half = 1;
neighbor->requests[irequest]->full_cluster = 0;
}
/* ----------------------------------------------------------------------
neighbor callback to inform pair style of neighbor list to use
regular or rRESPA
------------------------------------------------------------------------- */
void PairLJCutKokkosLight::init_list(int id, NeighList *ptr)
{
if (id == 0) list = ptr;
else if (id == 1) listinner = ptr;
else if (id == 2) listmiddle = ptr;
else if (id == 3) listouter = ptr;
}
/* ----------------------------------------------------------------------
init for one type pair i,j and corresponding j,i
------------------------------------------------------------------------- */
double PairLJCutKokkosLight::init_one(int i, int j)
{
if (setflag[i][j] == 0) {
epsilon(i,j) = mix_energy(epsilon(i,i),epsilon(j,j),
sigma(i,i),sigma(j,j));
sigma(i,j) = mix_distance(sigma(i,i),sigma(j,j));
cut(i,j) = mix_distance(cut(i,i),cut(j,j));
}
lj1(i,j) = 48.0 * epsilon(i,j) * pow(sigma(i,j),12.0);
lj2(i,j) = 24.0 * epsilon(i,j) * pow(sigma(i,j),6.0);
lj3(i,j) = 4.0 * epsilon(i,j) * pow(sigma(i,j),12.0);
lj4(i,j) = 4.0 * epsilon(i,j) * pow(sigma(i,j),6.0);
if (offset_flag) {
double ratio = sigma(i,j) / cut(i,j);
offset(i,j) = 4.0 * epsilon(i,j) * (pow(ratio,12.0) - pow(ratio,6.0));
} else offset(i,j) = 0.0;
lj1(j,i) = lj1(i,j);
lj2(j,i) = lj2(i,j);
lj3(j,i) = lj3(i,j);
lj4(j,i) = lj4(i,j);
offset(j,i) = offset(i,j);
// check interior rRESPA cutoff
if (cut_respa && cut(i,j) < cut_respa[3])
error->all(FLERR,"Pair cutoff < Respa interior cutoff");
// compute I,J contribution to long-range tail correction
// count total # of atoms of type I and J via Allreduce
if (tail_flag) {
int *type = atom->type;
int nlocal = atom->nlocal;
double count[2],all[2];
count[0] = count[1] = 0.0;
for (int k = 0; k < nlocal; k++) {
if (type[k] == i) count[0] += 1.0;
if (type[k] == j) count[1] += 1.0;
}
MPI_Allreduce(count,all,2,MPI_DOUBLE,MPI_SUM,world);
double sig2 = sigma(i,j)*sigma(i,j);
double sig6 = sig2*sig2*sig2;
double rc3 = cut(i,j)*cut(i,j)*cut(i,j);
double rc6 = rc3*rc3;
double rc9 = rc3*rc6;
etail_ij = 8.0*MY_PI*all[0]*all[1]*epsilon(i,j) *
sig6 * (sig6 - 3.0*rc6) / (9.0*rc9);
ptail_ij = 16.0*MY_PI*all[0]*all[1]*epsilon(i,j) *
sig6 * (2.0*sig6 - 3.0*rc6) / (9.0*rc9);
}
return cut(i,j);
}
/* ----------------------------------------------------------------------
proc 0 writes to restart file
------------------------------------------------------------------------- */
void PairLJCutKokkosLight::write_restart(FILE *fp)
{
write_restart_settings(fp);
int i,j;
for (i = 1; i <= atom->ntypes; i++)
for (j = i; j <= atom->ntypes; j++) {
fwrite(&setflag[i][j],sizeof(int),1,fp);
if (setflag[i][j]) {
fwrite(&epsilon(i,j),sizeof(double),1,fp);
fwrite(&sigma(i,j),sizeof(double),1,fp);
fwrite(&cut(i,j),sizeof(double),1,fp);
}
}
}
/* ----------------------------------------------------------------------
proc 0 reads from restart file, bcasts
------------------------------------------------------------------------- */
void PairLJCutKokkosLight::read_restart(FILE *fp)
{
read_restart_settings(fp);
allocate();
int i,j;
int me = comm->me;
for (i = 1; i <= atom->ntypes; i++)
for (j = i; j <= atom->ntypes; j++) {
if (me == 0) fread(&setflag[i][j],sizeof(int),1,fp);
MPI_Bcast(&setflag[i][j],1,MPI_INT,0,world);
if (setflag[i][j]) {
if (me == 0) {
fread(&epsilon(i,j),sizeof(double),1,fp);
fread(&sigma(i,j),sizeof(double),1,fp);
fread(&cut(i,j),sizeof(double),1,fp);
}
MPI_Bcast(&epsilon(i,j),1,MPI_DOUBLE,0,world);
MPI_Bcast(&sigma(i,j),1,MPI_DOUBLE,0,world);
MPI_Bcast(&cut(i,j),1,MPI_DOUBLE,0,world);
}
}
}
/* ----------------------------------------------------------------------
proc 0 writes to restart file
------------------------------------------------------------------------- */
void PairLJCutKokkosLight::write_restart_settings(FILE *fp)
{
fwrite(&cut_global,sizeof(double),1,fp);
fwrite(&offset_flag,sizeof(int),1,fp);
fwrite(&mix_flag,sizeof(int),1,fp);
fwrite(&tail_flag,sizeof(int),1,fp);
}
/* ----------------------------------------------------------------------
proc 0 reads from restart file, bcasts
------------------------------------------------------------------------- */
void PairLJCutKokkosLight::read_restart_settings(FILE *fp)
{
int me = comm->me;
if (me == 0) {
fread(&cut_global,sizeof(double),1,fp);
fread(&offset_flag,sizeof(int),1,fp);
fread(&mix_flag,sizeof(int),1,fp);
fread(&tail_flag,sizeof(int),1,fp);
}
MPI_Bcast(&cut_global,1,MPI_DOUBLE,0,world);
MPI_Bcast(&offset_flag,1,MPI_INT,0,world);
MPI_Bcast(&mix_flag,1,MPI_INT,0,world);
MPI_Bcast(&tail_flag,1,MPI_INT,0,world);
}
/* ----------------------------------------------------------------------
proc 0 writes to data file
------------------------------------------------------------------------- */
void PairLJCutKokkosLight::write_data(FILE *fp)
{
for (int i = 1; i <= atom->ntypes; i++)
fprintf(fp,"%d %g %g\n",i,epsilon(i,i),sigma(i,i));
}
/* ----------------------------------------------------------------------
proc 0 writes all pairs to data file
------------------------------------------------------------------------- */
void PairLJCutKokkosLight::write_data_all(FILE *fp)
{
for (int i = 1; i <= atom->ntypes; i++)
for (int j = i; j <= atom->ntypes; j++)
fprintf(fp,"%d %d %g %g %g\n",i,j,epsilon(i,j),sigma(i,j),cut(i,j));
}
/* ---------------------------------------------------------------------- */
double PairLJCutKokkosLight::single(int i, int j, int itype, int jtype, double rsq,
double factor_coul, double factor_lj,
double &fforce)
{
double r2inv,r6inv,forcelj,philj;
r2inv = 1.0/rsq;
r6inv = r2inv*r2inv*r2inv;
forcelj = r6inv * (lj1(itype,jtype)*r6inv - lj2(itype,jtype));
fforce = factor_lj*forcelj*r2inv;
philj = r6inv*(lj3(itype,jtype)*r6inv-lj4(itype,jtype)) -
offset(itype,jtype);
return factor_lj*philj;
}
/* ---------------------------------------------------------------------- */
void *PairLJCutKokkosLight::extract(const char *str, int &dim)
{
dim = 2;
if (strcmp(str,"epsilon") == 0) return (void *) epsilon.ptr_on_device();
if (strcmp(str,"sigma") == 0) return (void *) sigma.ptr_on_device();
return NULL;
}
c4science · Help