ICODE-ADC/lammps
4
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Updates to the core-shell gpu pair styles

Browse Source
This commit is contained in:
Trung Nguyen
2018-06-20 10:43:13 -05:00
parent 3ae8fdccd8
commit a9b794b2ab
24 changed files with 454 additions and 299 deletions
Download Patch File Download Diff File Expand all files Collapse all files

26
lib/gpu/Nvidia.makefile
View File

@ -81,7 +81,8 @@ OBJS = $(OBJ_DIR)/lal_atom.o $(OBJ_DIR)/lal_ans.o \
$(OBJ_DIR)/lal_dipole_long_lj.o $(OBJ_DIR)/lal_dipole_long_lj_ext.o \
$(OBJ_DIR)/lal_lj_expand_coul_long.o $(OBJ_DIR)/lal_lj_expand_coul_long_ext.o \
$(OBJ_DIR)/lal_coul_long_cs.o $(OBJ_DIR)/lal_coul_long_cs_ext.o \
$(OBJ_DIR)/lal_born_coul_long_cs.o $(OBJ_DIR)/lal_born_coul_long_cs_ext.o
$(OBJ_DIR)/lal_born_coul_long_cs.o $(OBJ_DIR)/lal_born_coul_long_cs_ext.o \
$(OBJ_DIR)/lal_born_coul_wolf_cs.o $(OBJ_DIR)/lal_born_coul_wolf_cs_ext.o
CBNS = $(OBJ_DIR)/device.cubin $(OBJ_DIR)/device_cubin.h \
$(OBJ_DIR)/atom.cubin $(OBJ_DIR)/atom_cubin.h \
@ -141,7 +142,8 @@ CBNS = $(OBJ_DIR)/device.cubin $(OBJ_DIR)/device_cubin.h \
$(OBJ_DIR)/dipole_long_lj.cubin $(OBJ_DIR)/dipole_long_lj_cubin.h \
$(OBJ_DIR)/lj_expand_coul_long.cubin $(OBJ_DIR)/lj_expand_coul_long_cubin.h \
$(OBJ_DIR)/coul_long_cs.cubin $(OBJ_DIR)/coul_long_cs_cubin.h \
$(OBJ_DIR)/born_coul_long_cs.cubin $(OBJ_DIR)/born_coul_long_cs_cubin.h
$(OBJ_DIR)/born_coul_long_cs.cubin $(OBJ_DIR)/born_coul_long_cs_cubin.h \
$(OBJ_DIR)/born_coul_wolf_cs.cubin $(OBJ_DIR)/born_coul_wolf_cs_cubin.h
all: $(OBJ_DIR) $(GPU_LIB) $(EXECS)
@ -847,10 +849,10 @@ $(OBJ_DIR)/coul_long_cs.cubin: lal_coul_long_cs.cu lal_precision.h lal_preproces
$(OBJ_DIR)/coul_long_cs_cubin.h: $(OBJ_DIR)/coul_long_cs.cubin $(OBJ_DIR)/coul_long_cs.cubin
$(BIN2C) -c -n coul_long_cs $(OBJ_DIR)/coul_long_cs.cubin > $(OBJ_DIR)/coul_long_cs_cubin.h
$(OBJ_DIR)/lal_coul_long_cs.o: $(ALL_H) lal_coul_long_cs.h lal_coul_long_cs.cpp $(OBJ_DIR)/coul_long_cs_cubin.h $(OBJ_DIR)/lal_base_charge.o
$(OBJ_DIR)/lal_coul_long_cs.o: $(ALL_H) lal_coul_long_cs.h lal_coul_long_cs.cpp $(OBJ_DIR)/coul_long_cs_cubin.h $(OBJ_DIR)/lal_base_charge.o $(OBJ_DIR)/lal_coul_long.o
$(CUDR) -o $@ -c lal_coul_long_cs.cpp -I$(OBJ_DIR)
$(OBJ_DIR)/lal_coul_long_cs_ext.o: $(ALL_H) lal_coul_long_cs.h lal_coul_long_cs_ext.cpp lal_base_charge.h
$(OBJ_DIR)/lal_coul_long_cs_ext.o: $(ALL_H) lal_coul_long_cs.h lal_coul_long_cs_ext.cpp lal_coul_long.h
$(CUDR) -o $@ -c lal_coul_long_cs_ext.cpp -I$(OBJ_DIR)
$(OBJ_DIR)/born_coul_long_cs.cubin: lal_born_coul_long_cs.cu lal_precision.h lal_preprocessor.h
@ -859,12 +861,24 @@ $(OBJ_DIR)/born_coul_long_cs.cubin: lal_born_coul_long_cs.cu lal_precision.h lal
$(OBJ_DIR)/born_coul_long_cs_cubin.h: $(OBJ_DIR)/born_coul_long_cs.cubin $(OBJ_DIR)/born_coul_long_cs.cubin
$(BIN2C) -c -n born_coul_long_cs $(OBJ_DIR)/born_coul_long_cs.cubin > $(OBJ_DIR)/born_coul_long_cs_cubin.h
$(OBJ_DIR)/lal_born_coul_long_cs.o: $(ALL_H) lal_born_coul_long_cs.h lal_born_coul_long_cs.cpp $(OBJ_DIR)/born_coul_long_cs_cubin.h $(OBJ_DIR)/lal_base_charge.o
$(OBJ_DIR)/lal_born_coul_long_cs.o: $(ALL_H) lal_born_coul_long_cs.h lal_born_coul_long_cs.cpp $(OBJ_DIR)/born_coul_long_cs_cubin.h $(OBJ_DIR)/lal_base_charge.o $(OBJ_DIR)/lal_born_coul_long.o
$(CUDR) -o $@ -c lal_born_coul_long_cs.cpp -I$(OBJ_DIR)
$(OBJ_DIR)/lal_born_coul_long_cs_ext.o: $(ALL_H) lal_born_coul_long_cs.h lal_born_coul_long_cs_ext.cpp lal_base_charge.h
$(OBJ_DIR)/lal_born_coul_long_cs_ext.o: $(ALL_H) lal_born_coul_long_cs.h lal_born_coul_long_cs_ext.cpp lal_born_coul_long.h
$(CUDR) -o $@ -c lal_born_coul_long_cs_ext.cpp -I$(OBJ_DIR)
$(OBJ_DIR)/born_coul_wolf_cs.cubin: lal_born_coul_wolf_cs.cu lal_precision.h lal_preprocessor.h
$(CUDA) --cubin -DNV_KERNEL -o $@ lal_born_coul_wolf_cs.cu
$(OBJ_DIR)/born_coul_wolf_cs_cubin.h: $(OBJ_DIR)/born_coul_wolf_cs.cubin $(OBJ_DIR)/born_coul_wolf_cs.cubin
$(BIN2C) -c -n born_coul_wolf_cs $(OBJ_DIR)/born_coul_wolf_cs.cubin > $(OBJ_DIR)/born_coul_wolf_cs_cubin.h
$(OBJ_DIR)/lal_born_coul_wolf_cs.o: $(ALL_H) lal_born_coul_wolf_cs.h lal_born_coul_wolf_cs.cpp $(OBJ_DIR)/born_coul_wolf_cs_cubin.h $(OBJ_DIR)/lal_base_charge.o $(OBJ_DIR)/lal_born_coul_wolf.o
$(CUDR) -o $@ -c lal_born_coul_wolf_cs.cpp -I$(OBJ_DIR)
$(OBJ_DIR)/lal_born_coul_wolf_cs_ext.o: $(ALL_H) lal_born_coul_wolf_cs.h lal_born_coul_wolf_cs_ext.cpp lal_born_coul_wolf.h
$(CUDR) -o $@ -c lal_born_coul_wolf_cs_ext.cpp -I$(OBJ_DIR)
$(BIN_DIR)/nvc_get_devices: ./geryon/ucl_get_devices.cpp $(NVD_H)
$(CUDR) -o $@ ./geryon/ucl_get_devices.cpp -DUCL_CUDADR $(CUDA_LIB) -lcuda

22
lib/gpu/Opencl.makefile
View File

@ -70,7 +70,8 @@ OBJS = $(OBJ_DIR)/lal_atom.o $(OBJ_DIR)/lal_answer.o \
$(OBJ_DIR)/lal_dipole_long_lj.o $(OBJ_DIR)/lal_dipole_long_lj_ext.o \
$(OBJ_DIR)/lal_lj_expand_coul_long.o $(OBJ_DIR)/lal_lj_expand_coul_long_ext.o \
$(OBJ_DIR)/lal_coul_long_cs.o $(OBJ_DIR)/lal_coul_long_cs_ext.o \
$(OBJ_DIR)/lal_born_coul_long_cs.o $(OBJ_DIR)/lal_born_coul_long_cs_ext.o
$(OBJ_DIR)/lal_born_coul_long_cs.o $(OBJ_DIR)/lal_born_coul_long_cs_ext.o \
$(OBJ_DIR)/lal_born_coul_wolf_cs.o $(OBJ_DIR)/lal_born_coul_wolf_cs_ext.o
KERS = $(OBJ_DIR)/device_cl.h $(OBJ_DIR)/atom_cl.h \
$(OBJ_DIR)/neighbor_cpu_cl.h $(OBJ_DIR)/pppm_cl.h \
@ -101,7 +102,7 @@ KERS = $(OBJ_DIR)/device_cl.h $(OBJ_DIR)/atom_cl.h \
$(OBJ_DIR)/lj_cubic_cl.h $(OBJ_DIR)/vashishta_cl.h \
$(OBJ_DIR)/ufm_cl.h $(OBJ_DIR)/dipole_long_lj_cl.h \
$(OBJ_DIR)/lj_expand_coul_long_cl.h $(OBJ_DIR)/coul_long_cs_cl.h \
$(OBJ_DIR)/born_coul_long_cs_cl.h
$(OBJ_DIR)/born_coul_long_cs_cl.h $(OBJ_DIR)/born_coul_wolf_cs_cl.h
OCL_EXECS = $(BIN_DIR)/ocl_get_devices
@ -615,21 +616,30 @@ $(OBJ_DIR)/lal_lj_expand_coul_long_ext.o: $(ALL_H) lal_lj_expand_coul_long.h lal
$(OBJ_DIR)/coul_long_cs_cl.h: lal_coul_long_cs.cu $(PRE1_H)
$(BSH) ./geryon/file_to_cstr.sh coul_long_cs $(PRE1_H) lal_coul_long_cs.cu $(OBJ_DIR)/coul_long_cs_cl.h;
$(OBJ_DIR)/lal_coul_long_cs.o: $(ALL_H) lal_coul_long_cs.h lal_coul_long_cs.cpp $(OBJ_DIR)/coul_long_cs_cl.h $(OBJ_DIR)/coul_long_cs_cl.h $(OBJ_DIR)/lal_base_charge.o
$(OBJ_DIR)/lal_coul_long_cs.o: $(ALL_H) lal_coul_long_cs.h lal_coul_long_cs.cpp $(OBJ_DIR)/coul_long_cs_cl.h $(OBJ_DIR)/coul_long_cs_cl.h $(OBJ_DIR)/lal_base_charge.o $(OBJ_DIR)/lal_coul_long.o
$(OCL) -o $@ -c lal_coul_long_cs.cpp -I$(OBJ_DIR)
$(OBJ_DIR)/lal_coul_long_cs_ext.o: $(ALL_H) lal_coul_long_cs.h lal_coul_long_cs_ext.cpp lal_base_charge.h
$(OBJ_DIR)/lal_coul_long_cs_ext.o: $(ALL_H) lal_coul_long_cs.h lal_coul_long_cs_ext.cpp lal_coul_long.h
$(OCL) -o $@ -c lal_coul_long_cs_ext.cpp -I$(OBJ_DIR)
$(OBJ_DIR)/born_coul_long_cs_cl.h: lal_born_coul_long_cs.cu $(PRE1_H)
$(BSH) ./geryon/file_to_cstr.sh born_coul_long_cs $(PRE1_H) lal_born_coul_long_cs.cu $(OBJ_DIR)/born_coul_long_cs_cl.h;
$(OBJ_DIR)/lal_born_coul_long_cs.o: $(ALL_H) lal_born_coul_long_cs.h lal_born_coul_long_cs.cpp $(OBJ_DIR)/born_coul_long_cs_cl.h $(OBJ_DIR)/born_coul_long_cs_cl.h $(OBJ_DIR)/lal_base_charge.o
$(OBJ_DIR)/lal_born_coul_long_cs.o: $(ALL_H) lal_born_coul_long_cs.h lal_born_coul_long_cs.cpp $(OBJ_DIR)/born_coul_long_cs_cl.h $(OBJ_DIR)/born_coul_long_cs_cl.h $(OBJ_DIR)/lal_base_charge.o $(OBJ_DIR)/lal_born_coul_long.o
$(OCL) -o $@ -c lal_born_coul_long_cs.cpp -I$(OBJ_DIR)
$(OBJ_DIR)/lal_born_coul_long_cs_ext.o: $(ALL_H) lal_born_coul_long_cs.h lal_born_coul_long_cs_ext.cpp lal_base_charge.h
$(OBJ_DIR)/lal_born_coul_long_cs_ext.o: $(ALL_H) lal_born_coul_long_cs.h lal_born_coul_long_cs_ext.cpp lal_born_coul_long.h
$(OCL) -o $@ -c lal_born_coul_long_cs_ext.cpp -I$(OBJ_DIR)
$(OBJ_DIR)/born_coul_wolf_cs_cl.h: lal_born_coul_wolf_cs.cu $(PRE1_H)
$(BSH) ./geryon/file_to_cstr.sh born_coul_wolf_cs $(PRE1_H) lal_born_coul_wolf_cs.cu $(OBJ_DIR)/born_coul_wolf_cs_cl.h;
$(OBJ_DIR)/lal_born_coul_wolf_cs.o: $(ALL_H) lal_born_coul_wolf_cs.h lal_born_coul_wolf_cs.cpp $(OBJ_DIR)/born_coul_wolf_cs_cl.h $(OBJ_DIR)/born_coul_wolf_cs_cl.h $(OBJ_DIR)/lal_base_charge.o $(OBJ_DIR)/lal_born_coul_wolf.o
$(OCL) -o $@ -c lal_born_coul_wolf_cs.cpp -I$(OBJ_DIR)
$(OBJ_DIR)/lal_born_coul_wolf_cs_ext.o: $(ALL_H) lal_born_coul_wolf_cs.h lal_born_coul_wolf_cs_ext.cpp lal_born_coul_wolf.h
$(OCL) -o $@ -c lal_born_coul_wolf_cs_ext.cpp -I$(OBJ_DIR)
$(BIN_DIR)/ocl_get_devices: ./geryon/ucl_get_devices.cpp $(OCL_H)
$(OCL) -o $@ ./geryon/ucl_get_devices.cpp -DUCL_OPENCL $(OCL_LINK)

2
lib/gpu/lal_born_coul_long.h
View File

@ -78,7 +78,7 @@ class BornCoulLong : public BaseCharge<numtyp, acctyp> {
numtyp _cut_coulsq, _qqrd2e, _g_ewald;
private:
protected:
bool _allocated;
void loop(const bool _eflag, const bool _vflag);
};

116
lib/gpu/lal_born_coul_long_cs.cpp
View File

@ -10,7 +10,7 @@
__________________________________________________________________________
begin :
email : trung.nguyen@northwestern.edu
email : ndactrung@gmail.com
***************************************************************************/
#ifdef USE_OPENCL
@ -28,21 +28,6 @@ using namespace LAMMPS_AL;
extern Device<PRECISION,ACC_PRECISION> device;
template <class numtyp, class acctyp>
BornCoulLongCST::BornCoulLongCS() : BaseCharge<numtyp,acctyp>(),
_allocated(false) {
}
template <class numtyp, class acctyp>
BornCoulLongCST::~BornCoulLongCS() {
clear();
}
template <class numtyp, class acctyp>
int BornCoulLongCST::bytes_per_atom(const int max_nbors) const {
return this->bytes_per_atom_atomic(max_nbors);
}
template <class numtyp, class acctyp>
int BornCoulLongCST::init(const int ntypes, double **host_cutsq, double **host_rhoinv,
double **host_born1, double **host_born2, double **host_born3,
@ -63,13 +48,13 @@ int BornCoulLongCST::init(const int ntypes, double **host_cutsq, double **host_r
// If atom type constants fit in shared memory use fast kernel
int lj_types=ntypes;
shared_types=false;
this->shared_types=false;
int max_shared_types=this->device->max_shared_types();
if (lj_types<=max_shared_types && this->_block_size>=max_shared_types) {
lj_types=max_shared_types;
shared_types=true;
this->shared_types=true;
}
_lj_types=lj_types;
this->_lj_types=lj_types;
// Allocate a host write buffer for data initialization
UCL_H_Vec<numtyp> host_write(lj_types*lj_types*32,*(this->ucl_device),
@ -78,98 +63,33 @@ int BornCoulLongCST::init(const int ntypes, double **host_cutsq, double **host_r
for (int i=0; i<lj_types*lj_types; i++)
host_write[i]=0.0;
coeff1.alloc(lj_types*lj_types,*(this->ucl_device),UCL_READ_ONLY);
this->atom->type_pack4(ntypes,lj_types,coeff1,host_write,host_rhoinv,
this->coeff1.alloc(lj_types*lj_types,*(this->ucl_device),UCL_READ_ONLY);
this->atom->type_pack4(ntypes,lj_types,this->coeff1,host_write,host_rhoinv,
host_born1,host_born2,host_born3);
coeff2.alloc(lj_types*lj_types,*(this->ucl_device),UCL_READ_ONLY);
this->atom->type_pack4(ntypes,lj_types,coeff2,host_write,host_a,host_c,
this->coeff2.alloc(lj_types*lj_types,*(this->ucl_device),UCL_READ_ONLY);
this->atom->type_pack4(ntypes,lj_types,this->coeff2,host_write,host_a,host_c,
host_d,host_offset);
cutsq_sigma.alloc(lj_types*lj_types,*(this->ucl_device),UCL_READ_ONLY);
this->atom->type_pack4(ntypes,lj_types,cutsq_sigma,host_write,host_cutsq,
this->cutsq_sigma.alloc(lj_types*lj_types,*(this->ucl_device),UCL_READ_ONLY);
this->atom->type_pack4(ntypes,lj_types,this->cutsq_sigma,host_write,host_cutsq,
host_cut_ljsq,host_sigma);
sp_lj.alloc(8,*(this->ucl_device),UCL_READ_ONLY);
this->sp_lj.alloc(8,*(this->ucl_device),UCL_READ_ONLY);
for (int i=0; i<4; i++) {
host_write[i]=host_special_lj[i];
host_write[i+4]=host_special_coul[i];
}
ucl_copy(sp_lj,host_write,8,false);
ucl_copy(this->sp_lj,host_write,8,false);
_cut_coulsq=host_cut_coulsq;
_qqrd2e=qqrd2e;
_g_ewald=g_ewald;
this->_cut_coulsq=host_cut_coulsq;
this->_qqrd2e=qqrd2e;
this->_g_ewald=g_ewald;
_allocated=true;
this->_max_bytes=coeff1.row_bytes()+coeff2.row_bytes()
+cutsq_sigma.row_bytes()+sp_lj.row_bytes();
this->_allocated=true;
this->_max_bytes=this->coeff1.row_bytes()+this->coeff2.row_bytes()
+this->cutsq_sigma.row_bytes()+this->sp_lj.row_bytes();
return 0;
}
template <class numtyp, class acctyp>
void BornCoulLongCST::clear() {
if (!_allocated)
return;
_allocated=false;
coeff1.clear();
coeff2.clear();
cutsq_sigma.clear();
sp_lj.clear();
this->clear_atomic();
}
template <class numtyp, class acctyp>
double BornCoulLongCST::host_memory_usage() const {
return this->host_memory_usage_atomic()+sizeof(BornCoulLongCS<numtyp,acctyp>);
}
// ---------------------------------------------------------------------------
// Calculate energies, forces, and torques
// ---------------------------------------------------------------------------
template <class numtyp, class acctyp>
void BornCoulLongCST::loop(const bool _eflag, const bool _vflag) {
// Compute the block size and grid size to keep all cores busy
const int BX=this->block_size();
int eflag, vflag;
if (_eflag)
eflag=1;
else
eflag=0;
if (_vflag)
vflag=1;
else
vflag=0;
int GX=static_cast<int>(ceil(static_cast<double>(this->ans->inum())/
(BX/this->_threads_per_atom)));
int ainum=this->ans->inum();
int nbor_pitch=this->nbor->nbor_pitch();
this->time_pair.start();
if (shared_types) {
this->k_pair_fast.set_size(GX,BX);
this->k_pair_fast.run(&this->atom->x, &coeff1, &coeff2, &sp_lj,
&this->nbor->dev_nbor,
&this->_nbor_data->begin(),
&this->ans->force,
&this->ans->engv, &eflag, &vflag,
&ainum, &nbor_pitch, &this->atom->q,
&cutsq_sigma, &_cut_coulsq, &_qqrd2e,
&_g_ewald, &this->_threads_per_atom);
} else {
this->k_pair.set_size(GX,BX);
this->k_pair.run(&this->atom->x, &coeff1, &coeff2, &_lj_types, &sp_lj,
&this->nbor->dev_nbor, &this->_nbor_data->begin(),
&this->ans->force, &this->ans->engv,
&eflag, &vflag, &ainum,
&nbor_pitch, &this->atom->q,
&cutsq_sigma, &_cut_coulsq,
&_qqrd2e, &_g_ewald, &this->_threads_per_atom);
}
this->time_pair.stop();
}
template class BornCoulLongCS<PRECISION,ACC_PRECISION>;

2
lib/gpu/lal_born_coul_long_cs.cu
View File

@ -10,7 +10,7 @@
// __________________________________________________________________________
//
// begin : June 2018
// email : trung.nguyen@northwestern.edu
// email : ndactrung@gmail.com
// ***************************************************************************/
#ifdef NV_KERNEL

49
lib/gpu/lal_born_coul_long_cs.h
View File

@ -1,7 +1,7 @@
/***************************************************************************
born_coul_long.h
born_coul_long_cs.h
-------------------
Trung Dac Nguyen (ORNL)
Trung Dac Nguyen (Northwestern)
Class for acceleration of the born/coul/long/cs pair style.
@ -10,21 +10,21 @@
__________________________________________________________________________
begin :
email : trung.nguyen@northwestern.edu
email : ndactrung@gmail.com
***************************************************************************/
#ifndef LAL_BORN_COUL_LONG_CS_H
#define LAL_BORN_COUL_LONG_CS_H
#include "lal_base_charge.h"
#include "lal_born_coul_long.h"
namespace LAMMPS_AL {
template <class numtyp, class acctyp>
class BornCoulLongCS : public BaseCharge<numtyp, acctyp> {
class BornCoulLongCS : public BornCoulLong<numtyp, acctyp> {
public:
BornCoulLongCS();
~BornCoulLongCS();
BornCoulLongCS() {}
~BornCoulLongCS() {}
/// Clear any previous data and set up for a new LAMMPS run
/** \param max_nbors initial number of rows in the neighbor matrix
@ -46,41 +46,6 @@ class BornCoulLongCS : public BaseCharge<numtyp, acctyp> {
const double gpu_split, FILE *screen, double **host_cut_ljsq,
const double host_cut_coulsq, double *host_special_coul,
const double qqrd2e, const double g_ewald);
/// Clear all host and device data
/** \note This is called at the beginning of the init() routine **/
void clear();
/// Returns memory usage on device per atom
int bytes_per_atom(const int max_nbors) const;
/// Total host memory used by library for pair style
double host_memory_usage() const;
// --------------------------- TYPE DATA --------------------------
/// coeff1.x = rhoinv, coeff1.y = born1, coeff1.z = born2,
/// coeff1.w = born3
UCL_D_Vec<numtyp4> coeff1;
/// coeff2.x = a, coeff2.y = c, coeff2.z = d, coeff2.w = offset
UCL_D_Vec<numtyp4> coeff2;
/// cutsq_sigma.x = cutsq, cutsq_sigma.y = cutsq_lj,
/// cutsq_sigma.z = sigma
UCL_D_Vec<numtyp4> cutsq_sigma;
/// Special LJ values [0-3] and Special Coul values [4-7]
UCL_D_Vec<numtyp> sp_lj;
/// If atom type constants fit in shared memory, use fast kernels
bool shared_types;
/// Number of atom types
int _lj_types;
numtyp _cut_coulsq, _qqrd2e, _g_ewald;
private:
bool _allocated;
void loop(const bool _eflag, const bool _vflag);
};
}

2
lib/gpu/lal_born_coul_long_cs_ext.cpp
View File

@ -10,7 +10,7 @@
__________________________________________________________________________
begin :
email : trung.nguyen@northwestern.edu
email : ndactrung@gmail.com
***************************************************************************/
#include <iostream>

2
lib/gpu/lal_born_coul_wolf.cpp
View File

@ -57,7 +57,7 @@ int BornCoulWolfT::init(const int ntypes, double **host_cutsq, double **host_rho
const double alf, const double e_shift, const double f_shift) {
int success;
success=this->init_atomic(nlocal,nall,max_nbors,maxspecial,cell_size,gpu_split,
_screen,born_coul_wolf,"k_born_wolf");
_screen,born_coul_wolf,"k_born_coul_wolf");
if (success!=0)
return success;

4
lib/gpu/lal_born_coul_wolf.cu
View File

@ -31,7 +31,7 @@ texture<int2> q_tex;
#define MY_PIS (acctyp)1.77245385090551602729
__kernel void k_born_wolf(const __global numtyp4 *restrict x_,
__kernel void k_born_coul_wolf(const __global numtyp4 *restrict x_,
const __global numtyp4 *restrict coeff1,
const __global numtyp4 *restrict coeff2,
const int lj_types,
@ -165,7 +165,7 @@ __kernel void k_born_wolf(const __global numtyp4 *restrict x_,
} // if ii
}
__kernel void k_born_wolf_fast(const __global numtyp4 *restrict x_,
__kernel void k_born_coul_wolf_fast(const __global numtyp4 *restrict x_,
const __global numtyp4 *restrict coeff1_in,
const __global numtyp4 *restrict coeff2_in,
const __global numtyp *restrict sp_lj_in,

6
lib/gpu/lal_born_coul_wolf.h
View File

@ -13,8 +13,8 @@
email : nguyentd@ornl.gov
***************************************************************************/
#ifndef LAL_BORN_COUL_LONG_H
#define LAL_BORN_COUL_LONG_H
#ifndef LAL_BORN_COUL_WOLF_H
#define LAL_BORN_COUL_WOLF_H
#include "lal_base_charge.h"
@ -79,7 +79,7 @@ class BornCoulWolf : public BaseCharge<numtyp, acctyp> {
numtyp _cut_coulsq,_qqrd2e,_alf,_e_shift,_f_shift;
private:
protected:
bool _allocated;
void loop(const bool _eflag, const bool _vflag);
};

2
lib/gpu/lal_coul_long.h
View File

@ -72,7 +72,7 @@ class CoulLong : public BaseCharge<numtyp, acctyp> {
numtyp _cut_coulsq, _qqrd2e, _g_ewald;
private:
protected:
bool _allocated;
void loop(const bool _eflag, const bool _vflag);
};

106
lib/gpu/lal_coul_long_cs.cpp
View File

@ -10,7 +10,7 @@
__________________________________________________________________________
begin : June 2018
email : trung.nguyen@northwestern.edu
email : ndactrung@gmail.com
***************************************************************************/
#if defined(USE_OPENCL)
@ -28,20 +28,6 @@ using namespace LAMMPS_AL;
extern Device<PRECISION,ACC_PRECISION> pair_gpu_device;
template <class numtyp, class acctyp>
CoulLongCST::CoulLongCS() : BaseCharge<numtyp,acctyp>(), _allocated(false) {
}
template <class numtyp, class acctyp>
CoulLongCST::~CoulLongCS() {
clear();
}
template <class numtyp, class acctyp>
int CoulLongCST::bytes_per_atom(const int max_nbors) const {
return this->bytes_per_atom_atomic(max_nbors);
}
template <class numtyp, class acctyp>
int CoulLongCST::init(const int ntypes, double **host_scale,
const int nlocal, const int nall, const int max_nbors,
@ -56,13 +42,13 @@ int CoulLongCST::init(const int ntypes, double **host_scale,
return success;
int lj_types=ntypes;
shared_types=false;
this->shared_types=false;
int max_shared_types=this->device->max_shared_types();
if (lj_types<=max_shared_types && this->_block_size>=max_shared_types) {
lj_types=max_shared_types;
shared_types=true;
this->shared_types=true;
}
_lj_types=lj_types;
this->_lj_types=lj_types;
// Allocate a host write buffer for data initialization
UCL_H_Vec<numtyp> host_write(lj_types*lj_types*32,*(this->ucl_device),
@ -71,88 +57,22 @@ int CoulLongCST::init(const int ntypes, double **host_scale,
for (int i=0; i<lj_types*lj_types; i++)
host_write[i]=0.0;
scale.alloc(lj_types*lj_types,*(this->ucl_device),UCL_READ_ONLY);
this->atom->type_pack1(ntypes,lj_types,scale,host_write,host_scale);
this->scale.alloc(lj_types*lj_types,*(this->ucl_device),UCL_READ_ONLY);
this->atom->type_pack1(ntypes,lj_types,this->scale,host_write,host_scale);
sp_cl.alloc(4,*(this->ucl_device),UCL_READ_ONLY);
this->sp_cl.alloc(4,*(this->ucl_device),UCL_READ_ONLY);
for (int i=0; i<4; i++) {
host_write[i]=host_special_coul[i];
}
ucl_copy(sp_cl,host_write,4,false);
ucl_copy(this->sp_cl,host_write,4,false);
_cut_coulsq=host_cut_coulsq;
_qqrd2e=qqrd2e;
_g_ewald=g_ewald;
this->_cut_coulsq=host_cut_coulsq;
this->_qqrd2e=qqrd2e;
this->_g_ewald=g_ewald;
_allocated=true;
this->_max_bytes=scale.row_bytes()+sp_cl.row_bytes();
this->_allocated=true;
this->_max_bytes=this->scale.row_bytes()+this->sp_cl.row_bytes();
return 0;
}
template <class numtyp, class acctyp>
void CoulLongCST::reinit(const int ntypes, double **host_scale) {
UCL_H_Vec<numtyp> hscale(_lj_types*_lj_types,*(this->ucl_device),
UCL_WRITE_ONLY);
this->atom->type_pack1(ntypes,_lj_types,scale,hscale,host_scale);
}
template <class numtyp, class acctyp>
void CoulLongCST::clear() {
if (!_allocated)
return;
_allocated=false;
scale.clear();
sp_cl.clear();
this->clear_atomic();
}
template <class numtyp, class acctyp>
double CoulLongCST::host_memory_usage() const {
return this->host_memory_usage_atomic()+sizeof(CoulLongCS<numtyp,acctyp>);
}
// ---------------------------------------------------------------------------
// Calculate energies, forces, and torques
// ---------------------------------------------------------------------------
template <class numtyp, class acctyp>
void CoulLongCST::loop(const bool _eflag, const bool _vflag) {
// Compute the block size and grid size to keep all cores busy
const int BX=this->block_size();
int eflag, vflag;
if (_eflag)
eflag=1;
else
eflag=0;
if (_vflag)
vflag=1;
else
vflag=0;
int GX=static_cast<int>(ceil(static_cast<double>(this->ans->inum())/
(BX/this->_threads_per_atom)));
int ainum=this->ans->inum();
int nbor_pitch=this->nbor->nbor_pitch();
this->time_pair.start();
if (shared_types) {
this->k_pair_fast.set_size(GX,BX);
this->k_pair_fast.run(&this->atom->x, &scale, &sp_cl,
&this->nbor->dev_nbor, &this->_nbor_data->begin(),
&this->ans->force, &this->ans->engv,
&eflag, &vflag, &ainum, &nbor_pitch,
&this->atom->q, &_cut_coulsq, &_qqrd2e, &_g_ewald,
&this->_threads_per_atom);
} else {
this->k_pair.set_size(GX,BX);
this->k_pair.run(&this->atom->x, &scale, &_lj_types, &sp_cl,
&this->nbor->dev_nbor, &this->_nbor_data->begin(),
&this->ans->force, &this->ans->engv, &eflag, &vflag,
&ainum, &nbor_pitch, &this->atom->q, &_cut_coulsq,
&_qqrd2e, &_g_ewald, &this->_threads_per_atom);
}
this->time_pair.stop();
}
template class CoulLongCS<PRECISION,ACC_PRECISION>;

2
lib/gpu/lal_coul_long_cs.cu
View File

@ -10,7 +10,7 @@
// __________________________________________________________________________
//
// begin : June 2018
// email : trung.nguyen@northwestern.edu
// email : ndactrung@gmail.com
// ***************************************************************************/
#ifdef NV_KERNEL

42
lib/gpu/lal_coul_long_cs.h
View File

@ -10,21 +10,21 @@
__________________________________________________________________________
begin : June 2018
email : trung.nguyen@northwestern.edu
email : ndactrung@gmail.com
***************************************************************************/
#ifndef LAL_COUL_LONG_CS_H
#define LAL_COUL_LONG_CS_H
#include "lal_base_charge.h"
#include "lal_coul_long.h"
namespace LAMMPS_AL {
template <class numtyp, class acctyp>
class CoulLongCS : public BaseCharge<numtyp, acctyp> {
class CoulLongCS : public CoulLong<numtyp, acctyp> {
public:
CoulLongCS();
~CoulLongCS();
CoulLongCS() {}
~CoulLongCS() {}
/// Clear any previous data and set up for a new LAMMPS run
/** \param max_nbors initial number of rows in the neighbor matrix
@ -43,38 +43,6 @@ class CoulLongCS : public BaseCharge<numtyp, acctyp> {
const double gpu_split, FILE *screen,
const double host_cut_coulsq, double *host_special_coul,
const double qqrd2e, const double g_ewald);
/// Send updated coeffs from host to device (to be compatible with fix adapt)
void reinit(const int ntypes, double **scale);
/// Clear all host and device data
/** \note This is called at the beginning of the init() routine **/
void clear();
/// Returns memory usage on device per atom
int bytes_per_atom(const int max_nbors) const;
/// Total host memory used by library for pair style
double host_memory_usage() const;
// --------------------------- TYPE DATA --------------------------
/// scale
UCL_D_Vec<numtyp> scale;
/// Special Coul values [0-3]
UCL_D_Vec<numtyp> sp_cl;
/// If atom type constants fit in shared memory, use fast kernels
bool shared_types;
/// Number of atom types
int _lj_types;
numtyp _cut_coulsq, _qqrd2e, _g_ewald;
private:
bool _allocated;
void loop(const bool _eflag, const bool _vflag);
};
}

2
lib/gpu/lal_coul_long_cs_ext.cpp
View File

@ -10,7 +10,7 @@
__________________________________________________________________________
begin : June 2018
email : trung.nguyen@northwestern.edu
email : ndactrung@gmail.com
***************************************************************************/
#include <iostream>

2
lib/gpu/lal_lj_expand_coul_long.cpp
View File

@ -10,7 +10,7 @@
__________________________________________________________________________
begin :
email : trung.nguyen@northwestern.edu
email : ndactrung@gmail.com
***************************************************************************/
#if defined(USE_OPENCL)

2
lib/gpu/lal_lj_expand_coul_long.cu
View File

@ -10,7 +10,7 @@
// __________________________________________________________________________
//
// begin :
// email : trung.nguyen@northwestern.edu
// email : ndactrung@gmail.com
// ***************************************************************************/
#ifdef NV_KERNEL

2
lib/gpu/lal_lj_expand_coul_long.h
View File

@ -10,7 +10,7 @@
__________________________________________________________________________
begin :
email : trung.nguyen@northwestern.edu
email : ndactrung@gmail.com
***************************************************************************/
#ifndef LAL_LJ_EXPAND_COUL_LONG_H

4
lib/gpu/lal_lj_expand_coul_long_ext.cpp
View File

@ -1,5 +1,5 @@
/***************************************************************************
lj_coul_long_ext.cpp
lj_expand_coul_long_ext.cpp
------------------------
Trung Nguyen (Northwestern)
@ -10,7 +10,7 @@
__________________________________________________________________________
begin :
email : trung.nguyen@northwestern.edu
email : ndactrung@gmail.com
***************************************************************************/
#include <iostream>