From 6cebfc12b5c904dc7f03659a3a1fdb52a145cd41 Mon Sep 17 00:00:00 2001 From: sjplimp Date: Thu, 20 Oct 2011 14:36:25 +0000 Subject: [PATCH] git-svn-id: svn://svn.icms.temple.edu/lammps-ro/trunk@7144 f3b2605a-c512-4ea7-a41b-209d697bcdaa --- src/USER-CUDA/pair_sw_cuda.cpp | 209 ++++++++++++++++++++++ src/USER-CUDA/pair_sw_cuda.h | 66 +++++++ src/USER-CUDA/pair_tersoff_cuda.cpp | 206 ++++++++++++++++++++++ src/USER-CUDA/pair_tersoff_cuda.h | 66 +++++++ src/USER-CUDA/pair_tersoff_zbl_cuda.cpp | 220 ++++++++++++++++++++++++ src/USER-CUDA/pair_tersoff_zbl_cuda.h | 53 ++++++ 6 files changed, 820 insertions(+) create mode 100644 src/USER-CUDA/pair_sw_cuda.cpp create mode 100644 src/USER-CUDA/pair_sw_cuda.h create mode 100644 src/USER-CUDA/pair_tersoff_cuda.cpp create mode 100644 src/USER-CUDA/pair_tersoff_cuda.h create mode 100644 src/USER-CUDA/pair_tersoff_zbl_cuda.cpp create mode 100644 src/USER-CUDA/pair_tersoff_zbl_cuda.h diff --git a/src/USER-CUDA/pair_sw_cuda.cpp b/src/USER-CUDA/pair_sw_cuda.cpp new file mode 100644 index 0000000000..601ad7f2cf --- /dev/null +++ b/src/USER-CUDA/pair_sw_cuda.cpp @@ -0,0 +1,209 @@ +/* ---------------------------------------------------------------------- + LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator + + Original Version: + http://lammps.sandia.gov, Sandia National Laboratories + Steve Plimpton, sjplimp@sandia.gov + + See the README file in the top-level LAMMPS directory. + + ----------------------------------------------------------------------- + + USER-CUDA Package and associated modifications: + https://sourceforge.net/projects/lammpscuda/ + + Christian Trott, christian.trott@tu-ilmenau.de + Lars Winterfeld, lars.winterfeld@tu-ilmenau.de + Theoretical Physics II, University of Technology Ilmenau, Germany + + See the README file in the USER-CUDA directory. + + This software is distributed under the GNU General Public License. +------------------------------------------------------------------------- */ + +/* ---------------------------------------------------------------------- + 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 +#include +#include +#include +#include "pair_sw_cuda.h" +#include "cuda_data.h" +#include "atom.h" +#include "comm.h" +#include "force.h" +#include "neighbor.h" +#include "neigh_list.h" +#include "neigh_request.h" +#include "cuda_neigh_list.h" +#include "update.h" +#include "integrate.h" +#include "respa.h" +#include "memory.h" +#include "error.h" +#include "cuda.h" + +using namespace LAMMPS_NS; + + + + +/* ---------------------------------------------------------------------- */ + +PairSWCuda::PairSWCuda(LAMMPS *lmp) : PairSW(lmp) +{ + cuda = lmp->cuda; + if(cuda == NULL) + error->all(FLERR,"You cannot use a /cuda class, without activating 'cuda' acceleration. Provide '-c on' as command-line argument to LAMMPS.."); + + allocated2 = false; + params_f = NULL; + cuda->setSystemParams(); + cuda->shared_data.pair.cudable_force = 1; + cuda->shared_data.pair.override_block_per_atom = 0; + cuda->shared_data.pair.neighall = true; + init = false; +} + +/* ---------------------------------------------------------------------- + remember pointer to arrays in cuda shared data +------------------------------------------------------------------------- */ + +void PairSWCuda::allocate() +{ + if(! allocated) PairSW::allocate(); + if(! allocated2) + { + allocated2 = true; + cuda->shared_data.pair.cutsq = cutsq; + cuda->shared_data.pair.special_lj = force->special_lj; + cuda->shared_data.pair.special_coul = force->special_coul; + } +} + +/* ---------------------------------------------------------------------- */ + +void PairSWCuda::compute(int eflag, int vflag) +{ + if(!init) {Cuda_PairSWCuda_Init(&cuda->shared_data,params_f,map, &elem2param[0][0][0],nelements); init=true;} + if (eflag || vflag) ev_setup(eflag,vflag); + if(eflag) cuda->cu_eng_vdwl->upload(); + if(vflag) cuda->cu_virial->upload(); + + Cuda_PairSWCuda(& cuda->shared_data, & cuda_neigh_list->sneighlist, eflag, vflag, eflag_atom, vflag_atom);//,&elem2param[0][0][0],map + if(not cuda->shared_data.pair.collect_forces_later) + { + if(eflag) cuda->cu_eng_vdwl->download(); + if(vflag) cuda->cu_virial->download(); + } +} + +/* ---------------------------------------------------------------------- */ + +void PairSWCuda::settings(int narg, char **arg) +{ + PairSW::settings(narg, arg); +} + +/* ---------------------------------------------------------------------- */ + +void PairSWCuda::coeff(int narg, char **arg) +{ + PairSW::coeff(narg, arg); + allocate(); + params_f = (ParamSW_Float *) memory->srealloc(params_f,maxparam*sizeof(ParamSW_Float), + "pair:params_f"); + for(int i=0;ishared_data.pair.cut_global = cutmax; +} + +void PairSWCuda::init_style() +{ + MYDBG(printf("# CUDA PairSWCuda::init_style start\n"); ) + + int irequest; + + irequest = neighbor->request(this); + neighbor->requests[irequest]->full = 1; + neighbor->requests[irequest]->half = 0; + neighbor->requests[irequest]->cudable = 1; + neighbor->requests[irequest]->ghost = 1; + + + MYDBG(printf("# CUDA PairSWCuda::init_style end\n"); ) +} + +void PairSWCuda::init_list(int id, NeighList *ptr) +{ + MYDBG(printf("# CUDA PairSWCuda::init_list\n");) + PairSW::init_list(id, ptr); + // right now we can only handle verlet (id 0), not respa + if(id == 0) cuda_neigh_list = cuda->registerNeighborList(ptr); + // see Neighbor::init() for details on lammps lists' logic + MYDBG(printf("# CUDA PairSWCuda::init_list end\n");) + cu_params_f = (ParamSW_Float*) CudaWrapper_AllocCudaData(sizeof(ParamSW_Float)*maxparam); + CudaWrapper_UploadCudaData((void*) params_f,(void*) cu_params_f,sizeof(ParamSW_Float)*maxparam); + cu_elem2param = new cCudaData ((int*) elem2param, nelements,nelements,nelements); + cu_elem2param->upload(); + cu_map = new cCudaData ( map,atom->ntypes+1 ); + cu_map->upload(); +} + +void PairSWCuda::ev_setup(int eflag, int vflag) +{ + int maxeatomold=maxeatom; + PairSW::ev_setup(eflag,vflag); + + if (eflag_atom && atom->nmax > maxeatomold) + {delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );} + + if (vflag_atom && atom->nmax > maxeatomold) + {delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );} +} + + diff --git a/src/USER-CUDA/pair_sw_cuda.h b/src/USER-CUDA/pair_sw_cuda.h new file mode 100644 index 0000000000..be9ba9bb3d --- /dev/null +++ b/src/USER-CUDA/pair_sw_cuda.h @@ -0,0 +1,66 @@ +/* ---------------------------------------------------------------------- + LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator + + Original Version: + http://lammps.sandia.gov, Sandia National Laboratories + Steve Plimpton, sjplimp@sandia.gov + + See the README file in the top-level LAMMPS directory. + + ----------------------------------------------------------------------- + + USER-CUDA Package and associated modifications: + https://sourceforge.net/projects/lammpscuda/ + + Christian Trott, christian.trott@tu-ilmenau.de + Lars Winterfeld, lars.winterfeld@tu-ilmenau.de + Theoretical Physics II, University of Technology Ilmenau, Germany + + See the README file in the USER-CUDA directory. + + This software is distributed under the GNU General Public License. +------------------------------------------------------------------------- */ + +#ifdef PAIR_CLASS + +PairStyle(sw/cuda,PairSWCuda) + +#else + +#ifndef PAIR_SW_CUDA_H +#define PAIR_SW_CUDA_H + +#include "pair_sw_cuda_cu.h" +#include "pair_sw.h" +#include "cuda_data.h" + +namespace LAMMPS_NS { + +class PairSWCuda : public PairSW +{ + public: + PairSWCuda(class LAMMPS *); + void compute(int, int); + void settings(int, char **); + void coeff(int, char **); + void init_list(int, class NeighList *); + void init_style(); + void ev_setup(int eflag, int vflag); + protected: + + class Cuda *cuda; + void allocate(); + bool allocated2; + class CudaNeighList* cuda_neigh_list; + ParamSW_Float* params_f; + ParamSW_Float* cu_params_f; + cCudaData* cu_elem2param; + cCudaData* cu_map; + bool init; + bool iszbl; +}; + +} + +#endif +#endif diff --git a/src/USER-CUDA/pair_tersoff_cuda.cpp b/src/USER-CUDA/pair_tersoff_cuda.cpp new file mode 100644 index 0000000000..4f1dba4e31 --- /dev/null +++ b/src/USER-CUDA/pair_tersoff_cuda.cpp @@ -0,0 +1,206 @@ +/* ---------------------------------------------------------------------- + LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator + + Original Version: + http://lammps.sandia.gov, Sandia National Laboratories + Steve Plimpton, sjplimp@sandia.gov + + See the README file in the top-level LAMMPS directory. + + ----------------------------------------------------------------------- + + USER-CUDA Package and associated modifications: + https://sourceforge.net/projects/lammpscuda/ + + Christian Trott, christian.trott@tu-ilmenau.de + Lars Winterfeld, lars.winterfeld@tu-ilmenau.de + Theoretical Physics II, University of Technology Ilmenau, Germany + + See the README file in the USER-CUDA directory. + + This software is distributed under the GNU General Public License. +------------------------------------------------------------------------- */ + +/* ---------------------------------------------------------------------- + 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 +#include +#include +#include +#include "pair_tersoff_cuda.h" +#include "cuda_data.h" +#include "atom.h" +#include "comm.h" +#include "force.h" +#include "neighbor.h" +#include "neigh_list.h" +#include "neigh_request.h" +#include "cuda_neigh_list.h" +#include "update.h" +#include "integrate.h" +#include "respa.h" +#include "memory.h" +#include "error.h" +#include "cuda.h" + +using namespace LAMMPS_NS; + + + + +/* ---------------------------------------------------------------------- */ + +PairTersoffCuda::PairTersoffCuda(LAMMPS *lmp) : PairTersoff(lmp) +{ + cuda = lmp->cuda; + if(cuda == NULL) + error->all(FLERR,"You cannot use a /cuda class, without activating 'cuda' acceleration. Provide '-c on' as command-line argument to LAMMPS.."); + + allocated2 = false; + params_f = NULL; + cuda->setSystemParams(); + cuda->shared_data.pair.cudable_force = 1; + cuda->shared_data.pair.override_block_per_atom = 0; + cuda->shared_data.pair.neighall = true; + init = false; + iszbl = false; +} + +/* ---------------------------------------------------------------------- + remember pointer to arrays in cuda shared data +------------------------------------------------------------------------- */ + +void PairTersoffCuda::allocate() +{ + if(! allocated) PairTersoff::allocate(); + if(! allocated2) + { + allocated2 = true; + cuda->shared_data.pair.cutsq = cutsq; + cuda->shared_data.pair.special_lj = force->special_lj; + cuda->shared_data.pair.special_coul = force->special_coul; + } +} + +/* ---------------------------------------------------------------------- */ + +void PairTersoffCuda::compute(int eflag, int vflag) +{ + if(!init) {Cuda_PairTersoffCuda_Init(&cuda->shared_data,params_f,map, &elem2param[0][0][0],nelements,iszbl); init=true;} + if (eflag || vflag) ev_setup(eflag,vflag); + if(eflag) cuda->cu_eng_vdwl->upload(); + if(vflag) cuda->cu_virial->upload(); + + Cuda_PairTersoffCuda(& cuda->shared_data, & cuda_neigh_list->sneighlist, eflag, vflag, eflag_atom, vflag_atom);//,&elem2param[0][0][0],map + if(not cuda->shared_data.pair.collect_forces_later) + { + if(eflag) cuda->cu_eng_vdwl->download(); + if(vflag) cuda->cu_virial->download(); + } +} + +/* ---------------------------------------------------------------------- */ + +void PairTersoffCuda::settings(int narg, char **arg) +{ + PairTersoff::settings(narg, arg); +} + +/* ---------------------------------------------------------------------- */ + +void PairTersoffCuda::coeff(int narg, char **arg) +{ + PairTersoff::coeff(narg, arg); + allocate(); + params_f = (Param_Float *) memory->srealloc(params_f,maxparam*sizeof(Param_Float), + "pair:params_f"); + for(int i=0;ishared_data.pair.cut_global = cutmax; +} + +void PairTersoffCuda::init_style() +{ + MYDBG(printf("# CUDA PairTersoffCuda::init_style start\n"); ) + + int irequest; + + irequest = neighbor->request(this); + neighbor->requests[irequest]->full = 1; + neighbor->requests[irequest]->half = 0; + neighbor->requests[irequest]->cudable = 1; + neighbor->requests[irequest]->ghost = 1; + + + MYDBG(printf("# CUDA PairTersoffCuda::init_style end\n"); ) +} + +void PairTersoffCuda::init_list(int id, NeighList *ptr) +{ + MYDBG(printf("# CUDA PairTersoffCuda::init_list\n");) + PairTersoff::init_list(id, ptr); + // right now we can only handle verlet (id 0), not respa + if(id == 0) cuda_neigh_list = cuda->registerNeighborList(ptr); + // see Neighbor::init() for details on lammps lists' logic + MYDBG(printf("# CUDA PairTersoffCuda::init_list end\n");) + cu_params_f = (Param_Float*) CudaWrapper_AllocCudaData(sizeof(Param_Float)*maxparam); + CudaWrapper_UploadCudaData((void*) params_f,(void*) cu_params_f,sizeof(Param_Float)*maxparam); + cu_elem2param = new cCudaData ((int*) elem2param, nelements,nelements,nelements); + cu_elem2param->upload(); + cu_map = new cCudaData ( map,atom->ntypes+1 ); + cu_map->upload(); +} + +void PairTersoffCuda::ev_setup(int eflag, int vflag) +{ + int maxeatomold=maxeatom; + PairTersoff::ev_setup(eflag,vflag); + + if (eflag_atom && atom->nmax > maxeatomold) + {delete cuda->cu_eatom; cuda->cu_eatom = new cCudaData ((double*)eatom, & cuda->shared_data.atom.eatom , atom->nmax );} + + if (vflag_atom && atom->nmax > maxeatomold) + {delete cuda->cu_vatom; cuda->cu_vatom = new cCudaData ((double*)vatom, & cuda->shared_data.atom.vatom , atom->nmax, 6 );} +} + + diff --git a/src/USER-CUDA/pair_tersoff_cuda.h b/src/USER-CUDA/pair_tersoff_cuda.h new file mode 100644 index 0000000000..34a06f91a9 --- /dev/null +++ b/src/USER-CUDA/pair_tersoff_cuda.h @@ -0,0 +1,66 @@ +/* ---------------------------------------------------------------------- + LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator + + Original Version: + http://lammps.sandia.gov, Sandia National Laboratories + Steve Plimpton, sjplimp@sandia.gov + + See the README file in the top-level LAMMPS directory. + + ----------------------------------------------------------------------- + + USER-CUDA Package and associated modifications: + https://sourceforge.net/projects/lammpscuda/ + + Christian Trott, christian.trott@tu-ilmenau.de + Lars Winterfeld, lars.winterfeld@tu-ilmenau.de + Theoretical Physics II, University of Technology Ilmenau, Germany + + See the README file in the USER-CUDA directory. + + This software is distributed under the GNU General Public License. +------------------------------------------------------------------------- */ + +#ifdef PAIR_CLASS + +PairStyle(tersoff/cuda,PairTersoffCuda) + +#else + +#ifndef PAIR_TERSOFF_CUDA_H +#define PAIR_TERSOFF_CUDA_H + +#include "pair_tersoff_cuda_cu.h" +#include "pair_tersoff.h" +#include "cuda_data.h" + +namespace LAMMPS_NS { + +class PairTersoffCuda : public PairTersoff +{ + public: + PairTersoffCuda(class LAMMPS *); + void compute(int, int); + void settings(int, char **); + void coeff(int, char **); + void init_list(int, class NeighList *); + void init_style(); + void ev_setup(int eflag, int vflag); + protected: + + class Cuda *cuda; + void allocate(); + bool allocated2; + class CudaNeighList* cuda_neigh_list; + Param_Float* params_f; + Param_Float* cu_params_f; + cCudaData* cu_elem2param; + cCudaData* cu_map; + bool init; + bool iszbl; +}; + +} + +#endif +#endif diff --git a/src/USER-CUDA/pair_tersoff_zbl_cuda.cpp b/src/USER-CUDA/pair_tersoff_zbl_cuda.cpp new file mode 100644 index 0000000000..45236da515 --- /dev/null +++ b/src/USER-CUDA/pair_tersoff_zbl_cuda.cpp @@ -0,0 +1,220 @@ +/* ---------------------------------------------------------------------- + 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: Aidan Thompson (SNL) - original Tersoff implementation + David Farrell (NWU) - ZBL addition +------------------------------------------------------------------------- */ + +#include "math.h" +#include "stdio.h" +#include "stdlib.h" +#include "string.h" +#include "pair_tersoff_zbl_cuda.h" +#include "atom.h" +#include "update.h" +#include "neighbor.h" +#include "neigh_list.h" +#include "neigh_request.h" +#include "force.h" +#include "comm.h" +#include "memory.h" +#include "error.h" +#include "math_const.h" + +using namespace LAMMPS_NS; +using namespace MathConst; + +#define MAXLINE 1024 +#define DELTA 4 + +/* ---------------------------------------------------------------------- */ + +PairTersoffZBLCuda::PairTersoffZBLCuda(LAMMPS *lmp) : PairTersoffCuda(lmp) +{ + // hard-wired constants in metal or real units + // a0 = Bohr radius + // epsilon0 = permittivity of vacuum = q / energy-distance units + // e = unit charge + // 1 Kcal/mole = 0.043365121 eV + + if (strcmp(update->unit_style,"metal") == 0) { + global_a_0 = 0.529; + global_epsilon_0 = 0.00552635; + global_e = 1.0; + } else if (strcmp(update->unit_style,"real") == 0) { + global_a_0 = 0.529; + global_epsilon_0 = 0.00552635 * 0.043365121; + global_e = 1.0; + } else error->all(FLERR,"Pair tersoff/zbl requires metal or real units"); + iszbl = true; +} + +/* ---------------------------------------------------------------------- */ + +void PairTersoffZBLCuda::read_file(char *file) +{ + int params_per_line = 21; + char **words = new char*[params_per_line+1]; + + delete [] params; + params = NULL; + nparams = 0; + + // open file on proc 0 + + FILE *fp; + if (comm->me == 0) { + fp = fopen(file,"r"); + if (fp == NULL) { + char str[128]; + sprintf(str,"Cannot open Tersoff potential file %s",file); + error->one(FLERR,str); + } + } + + // read each line out of file, skipping blank lines or leading '#' + // store line of params if all 3 element tags are in element list + + int n,nwords,ielement,jelement,kelement; + char line[MAXLINE],*ptr; + int eof = 0; + + while (1) { + if (comm->me == 0) { + ptr = fgets(line,MAXLINE,fp); + if (ptr == NULL) { + eof = 1; + fclose(fp); + } else n = strlen(line) + 1; + } + MPI_Bcast(&eof,1,MPI_INT,0,world); + if (eof) break; + MPI_Bcast(&n,1,MPI_INT,0,world); + MPI_Bcast(line,n,MPI_CHAR,0,world); + + // strip comment, skip line if blank + + if (ptr = strchr(line,'#')) *ptr = '\0'; + nwords = atom->count_words(line); + if (nwords == 0) continue; + + // concatenate additional lines until have params_per_line words + + while (nwords < params_per_line) { + n = strlen(line); + if (comm->me == 0) { + ptr = fgets(&line[n],MAXLINE-n,fp); + if (ptr == NULL) { + eof = 1; + fclose(fp); + } else n = strlen(line) + 1; + } + MPI_Bcast(&eof,1,MPI_INT,0,world); + if (eof) break; + MPI_Bcast(&n,1,MPI_INT,0,world); + MPI_Bcast(line,n,MPI_CHAR,0,world); + if (ptr = strchr(line,'#')) *ptr = '\0'; + nwords = atom->count_words(line); + } + + if (nwords != params_per_line) + error->all(FLERR,"Incorrect format in Tersoff potential file"); + + // words = ptrs to all words in line + + nwords = 0; + words[nwords++] = strtok(line," \t\n\r\f"); + while (words[nwords++] = strtok(NULL," \t\n\r\f")) continue; + + // ielement,jelement,kelement = 1st args + // if all 3 args are in element list, then parse this line + // else skip to next line + + for (ielement = 0; ielement < nelements; ielement++) + if (strcmp(words[0],elements[ielement]) == 0) break; + if (ielement == nelements) continue; + for (jelement = 0; jelement < nelements; jelement++) + if (strcmp(words[1],elements[jelement]) == 0) break; + if (jelement == nelements) continue; + for (kelement = 0; kelement < nelements; kelement++) + if (strcmp(words[2],elements[kelement]) == 0) break; + if (kelement == nelements) continue; + + // load up parameter settings and error check their values + + if (nparams == maxparam) { + maxparam += DELTA; + params = (Param *) memory->srealloc(params,maxparam*sizeof(Param), + "pair:params"); + } + + params[nparams].ielement = ielement; + params[nparams].jelement = jelement; + params[nparams].kelement = kelement; + params[nparams].powerm = atof(words[3]); + params[nparams].gamma = atof(words[4]); + params[nparams].lam3 = atof(words[5]); + params[nparams].c = atof(words[6]); + params[nparams].d = atof(words[7]); + params[nparams].h = atof(words[8]); + params[nparams].powern = atof(words[9]); + params[nparams].beta = atof(words[10]); + params[nparams].lam2 = atof(words[11]); + params[nparams].bigb = atof(words[12]); + params[nparams].bigr = atof(words[13]); + params[nparams].bigd = atof(words[14]); + params[nparams].lam1 = atof(words[15]); + params[nparams].biga = atof(words[16]); + params[nparams].Z_i = atof(words[17]); + params[nparams].Z_j = atof(words[18]); + params[nparams].ZBLcut = atof(words[19]); + params[nparams].ZBLexpscale = atof(words[20]); + + // currently only allow m exponent of 1 or 3 + + params[nparams].powermint = int(params[nparams].powerm); + + if ( + params[nparams].lam3 < 0.0 || params[nparams].c < 0.0 || + params[nparams].d < 0.0 || params[nparams].powern < 0.0 || + params[nparams].beta < 0.0 || params[nparams].lam2 < 0.0 || + params[nparams].bigb < 0.0 || params[nparams].bigr < 0.0 || + params[nparams].bigd < 0.0 || + params[nparams].bigd > params[nparams].bigr || + params[nparams].lam3 < 0.0 || params[nparams].biga < 0.0 || + params[nparams].powerm - params[nparams].powermint != 0.0 || + (params[nparams].powermint != 3 && params[nparams].powermint != 1) || + params[nparams].gamma < 0.0 || + params[nparams].Z_i < 1.0 || params[nparams].Z_j < 1.0 || + params[nparams].ZBLcut < 0.0 || params[nparams].ZBLexpscale < 0.0) + error->all(FLERR,"Illegal Tersoff parameter"); + + nparams++; + } + + delete [] words; +} + +void PairTersoffZBLCuda::coeff(int narg, char **arg) +{ + PairTersoffCuda::coeff(narg, arg); + for(int i=0;i