Commit 6842a527 authored by Trung Nguyen's avatar Trung Nguyen
Browse files

Fixed bugs and added born/coul/long/cs/gpu and coul/long/cs/gpu, need tests

parent d094bf78

lib/gpu/Nvidia.makefile

+30 −2
Original line number Diff line number Diff line
@@ -79,7 +79,9 @@ OBJS = $(OBJ_DIR)/lal_atom.o $(OBJ_DIR)/lal_ans.o \ 
       $(OBJ_DIR)/lal_lj_cubic.o $(OBJ_DIR)/lal_lj_cubic_ext.o \

       $(OBJ_DIR)/lal_ufm.o $(OBJ_DIR)/lal_ufm_ext.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_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



CBNS = $(OBJ_DIR)/device.cubin $(OBJ_DIR)/device_cubin.h \

       $(OBJ_DIR)/atom.cubin $(OBJ_DIR)/atom_cubin.h \
@@ -137,7 +139,9 @@ CBNS = $(OBJ_DIR)/device.cubin $(OBJ_DIR)/device_cubin.h \ 
       $(OBJ_DIR)/lj_cubic.cubin $(OBJ_DIR)/lj_cubic_cubin.h \

       $(OBJ_DIR)/ufm.cubin $(OBJ_DIR)/ufm_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)/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



all: $(OBJ_DIR) $(GPU_LIB) $(EXECS)
@@ -837,6 +841,30 @@ $(OBJ_DIR)/lal_lj_expand_coul_long.o: $(ALL_H) lal_lj_expand_coul_long.h lal_lj_ 
$(OBJ_DIR)/lal_lj_expand_coul_long_ext.o: $(ALL_H) lal_lj_expand_coul_long.h lal_lj_expand_coul_long_ext.cpp lal_base_charge.h

	$(CUDR) -o $@ -c lal_lj_expand_coul_long_ext.cpp -I$(OBJ_DIR)



$(OBJ_DIR)/coul_long_cs.cubin: lal_coul_long_cs.cu lal_precision.h lal_preprocessor.h

	$(CUDA) --cubin -DNV_KERNEL -o $@ lal_coul_long_cs.cu



$(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

	$(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

	$(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

	$(CUDA) --cubin -DNV_KERNEL -o $@ lal_born_coul_long_cs.cu



$(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

	$(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

	$(CUDR) -o $@ -c lal_born_coul_long_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

lib/gpu/Opencl.makefile

+23 −2
Original line number Diff line number Diff line
@@ -68,7 +68,9 @@ OBJS = $(OBJ_DIR)/lal_atom.o $(OBJ_DIR)/lal_answer.o \ 
       $(OBJ_DIR)/lal_lj_cubic.o $(OBJ_DIR)/lal_lj_cubic_ext.o \

       $(OBJ_DIR)/lal_ufm.o $(OBJ_DIR)/lal_ufm_ext.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_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



KERS = $(OBJ_DIR)/device_cl.h $(OBJ_DIR)/atom_cl.h \

       $(OBJ_DIR)/neighbor_cpu_cl.h $(OBJ_DIR)/pppm_cl.h \
@@ -98,7 +100,8 @@ KERS = $(OBJ_DIR)/device_cl.h $(OBJ_DIR)/atom_cl.h \ 
       $(OBJ_DIR)/coul_debye_cl.h $(OBJ_DIR)/zbl_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)/lj_expand_coul_long_cl.h $(OBJ_DIR)/coul_long_cs_cl.h \

       $(OBJ_DIR)/born_coul_long_cs_cl.h





OCL_EXECS = $(BIN_DIR)/ocl_get_devices
@@ -609,6 +612,24 @@ $(OBJ_DIR)/lal_lj_expand_coul_long.o: $(ALL_H) lal_lj_expand_coul_long.h lal_lj_ 
$(OBJ_DIR)/lal_lj_expand_coul_long_ext.o: $(ALL_H) lal_lj_expand_coul_long.h lal_lj_expand_coul_long_ext.cpp lal_base_charge.h

	$(OCL) -o $@ -c lal_lj_expand_coul_long_ext.cpp -I$(OBJ_DIR)



$(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

	$(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

	$(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

	$(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

	$(OCL) -o $@ -c lal_born_coul_long_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)

lib/gpu/lal_born_coul_long.cpp

+1 −1
Original line number Diff line number Diff line
@@ -57,7 +57,7 @@ int BornCoulLongT::init(const int ntypes, double **host_cutsq, double **host_rho 
                       const double g_ewald) {

  int success;

  success=this->init_atomic(nlocal,nall,max_nbors,maxspecial,cell_size,gpu_split,

                            _screen,born_coul_long,"k_born_long");

                            _screen,born_coul_long,"k_born_coul_long");

  if (success!=0)

    return success;

lib/gpu/lal_born_coul_long.cu

+6 −6
Original line number Diff line number Diff line
@@ -29,7 +29,7 @@ texture<int2> q_tex; 
#define q_tex q_

#endif



__kernel void k_born_long(const __global numtyp4 *restrict x_,

__kernel void k_born_coul_long(const __global numtyp4 *restrict x_,

                          const __global numtyp4 *restrict coeff1,

                          const __global numtyp4 *restrict coeff2,

                          const int lj_types,
@@ -110,7 +110,7 @@ __kernel void k_born_long(const __global numtyp4 *restrict x_, 
          forcecoul = prefactor * (_erfc + EWALD_F*grij*expm2-factor_coul);

        } else forcecoul = (numtyp)0.0;



        if (rsq < cutsq_sigma[mtype].y) {

        if (rsq < cutsq_sigma[mtype].y) { // cut_ljsq

          numtyp r = ucl_sqrt(rsq);

          rexp = ucl_exp((cutsq_sigma[mtype].z-r)*coeff1[mtype].x);

          r6inv = r2inv*r2inv*r2inv;
@@ -127,7 +127,7 @@ __kernel void k_born_long(const __global numtyp4 *restrict x_, 
        if (eflag>0) {

          if (rsq < cut_coulsq)

            e_coul += prefactor*(_erfc-factor_coul);

          if (rsq < coeff1[mtype].w) {

          if (rsq < cutsq_sigma[mtype].y) {

            numtyp e=coeff2[mtype].x*rexp - coeff2[mtype].y*r6inv

              + coeff2[mtype].z*r2inv*r6inv;

            energy+=factor_lj*(e-coeff2[mtype].w);
@@ -149,7 +149,7 @@ __kernel void k_born_long(const __global numtyp4 *restrict x_, 
  } // if ii

}



__kernel void k_born_long_fast(const __global numtyp4 *restrict x_,

__kernel void k_born_coul_long_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,
@@ -232,7 +232,7 @@ __kernel void k_born_long_fast(const __global numtyp4 *restrict x_, 
          forcecoul = prefactor * (_erfc + EWALD_F*grij*expm2-factor_coul);

        } else forcecoul = (numtyp)0.0;



        if (rsq < cutsq_sigma[mtype].y) {

        if (rsq < cutsq_sigma[mtype].y) { // cut_ljsq

          numtyp r = ucl_sqrt(rsq);

          rexp = ucl_exp((cutsq_sigma[mtype].z-r)*coeff1[mtype].x);

          r6inv = r2inv*r2inv*r2inv;
@@ -249,7 +249,7 @@ __kernel void k_born_long_fast(const __global numtyp4 *restrict x_, 
        if (eflag>0) {

          if (rsq < cut_coulsq)

            e_coul += prefactor*(_erfc-factor_coul);

          if (rsq < coeff1[mtype].w) {

          if (rsq < cutsq_sigma[mtype].y) {

            numtyp e=coeff2[mtype].x*rexp - coeff2[mtype].y*r6inv

              + coeff2[mtype].z*r2inv*r6inv;

            energy+=factor_lj*(e-coeff2[mtype].w);

lib/gpu/lal_born_coul_long_cs.cpp

0 → 100644
+175 −0
Original line number Diff line number Diff line 
/***************************************************************************

                            born_coul_long_cs.cpp

                             -------------------

                           Trung Dac Nguyen (Northwestern)



  Class for acceleration of the born/coul/long/cs pair style.



 __________________________________________________________________________

    This file is part of the LAMMPS Accelerator Library (LAMMPS_AL)

 __________________________________________________________________________



    begin                :

    email                : trung.nguyen@northwestern.edu

 ***************************************************************************/



#ifdef USE_OPENCL

#include "born_coul_long_cs_cl.h"

#elif defined(USE_CUDART)

const char *born_coul_long_cs=0;

#else

#include "born_coul_long_cs_cubin.h"

#endif



#include "lal_born_coul_long_cs.h"

#include <cassert>

using namespace LAMMPS_AL;

#define BornCoulLongCST BornCoulLongCS<numtyp, acctyp>



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,

                       double **host_a, double **host_c, double **host_d,

                       double **host_sigma, double **host_offset,

                       double *host_special_lj, const int nlocal,

                       const int nall, const int max_nbors,

                       const int maxspecial, const double cell_size,

                       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) {

  int success;

  success=this->init_atomic(nlocal,nall,max_nbors,maxspecial,cell_size,gpu_split,

                            _screen,born_coul_long_cs,"k_born_coul_long_cs");

  if (success!=0)

    return success;



  // If atom type constants fit in shared memory use fast kernel

  int lj_types=ntypes;

  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;

  }

  _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),

                               UCL_WRITE_ONLY);



  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,

                         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,

                         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,

             host_cut_ljsq,host_sigma);



  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);



  _cut_coulsq=host_cut_coulsq;

  _qqrd2e=qqrd2e;

  _g_ewald=g_ewald;



  _allocated=true;

  this->_max_bytes=coeff1.row_bytes()+coeff2.row_bytes()

      +cutsq_sigma.row_bytes()+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>;

CRA Git | Maintained and supported by SUSTech CRA and CCSE