Commit cd435c0c authored by Axel Kohlmeyer's avatar Axel Kohlmeyer
Browse files

change references from cg_cmm to lj_sdk and from cmm to sdk

parent 548c589f

lib/gpu/Nvidia.makefile

+20 −20
Original line number Diff line number Diff line
@@ -43,8 +43,8 @@ OBJS = $(OBJ_DIR)/lal_atom.o $(OBJ_DIR)/lal_ans.o \ 
       $(OBJ_DIR)/lal_coul_long.o $(OBJ_DIR)/lal_coul_long_ext.o \

       $(OBJ_DIR)/lal_morse.o $(OBJ_DIR)/lal_morse_ext.o \

       $(OBJ_DIR)/lal_charmm_long.o $(OBJ_DIR)/lal_charmm_long_ext.o \

       $(OBJ_DIR)/lal_cg_cmm.o $(OBJ_DIR)/lal_cg_cmm_ext.o \

       $(OBJ_DIR)/lal_cg_cmm_long.o $(OBJ_DIR)/lal_cg_cmm_long_ext.o \

       $(OBJ_DIR)/lal_lj_sdk.o $(OBJ_DIR)/lal_lj_sdk_ext.o \

       $(OBJ_DIR)/lal_lj_sdk_long.o $(OBJ_DIR)/lal_lj_sdk_long_ext.o \

       $(OBJ_DIR)/lal_eam.o $(OBJ_DIR)/lal_eam_ext.o \

       $(OBJ_DIR)/lal_eam_fs_ext.o $(OBJ_DIR)/lal_eam_alloy_ext.o \

       $(OBJ_DIR)/lal_buck.o $(OBJ_DIR)/lal_buck_ext.o \
@@ -98,8 +98,8 @@ CBNS = $(OBJ_DIR)/device.cubin $(OBJ_DIR)/device_cubin.h \ 
       $(OBJ_DIR)/coul_long.cubin $(OBJ_DIR)/coul_long_cubin.h \

       $(OBJ_DIR)/morse.cubin $(OBJ_DIR)/morse_cubin.h \

       $(OBJ_DIR)/charmm_long.cubin $(OBJ_DIR)/charmm_long_cubin.h \

       $(OBJ_DIR)/cg_cmm.cubin $(OBJ_DIR)/cg_cmm_cubin.h \

       $(OBJ_DIR)/cg_cmm_long.cubin $(OBJ_DIR)/cg_cmm_long_cubin.h \

       $(OBJ_DIR)/lj_sdk.cubin $(OBJ_DIR)/lj_sdk_cubin.h \

       $(OBJ_DIR)/lj_sdk_long.cubin $(OBJ_DIR)/lj_sdk_long_cubin.h \

       $(OBJ_DIR)/eam.cubin $(OBJ_DIR)/eam_cubin.h \

       $(OBJ_DIR)/buck.cubin $(OBJ_DIR)/buck_cubin.h \

       $(OBJ_DIR)/buck_coul_long.cubin $(OBJ_DIR)/buck_coul_long_cubin.h \
@@ -391,29 +391,29 @@ $(OBJ_DIR)/lal_lj_expand.o: $(ALL_H) lal_lj_expand.h lal_lj_expand.cpp $(OBJ_DIR 
$(OBJ_DIR)/lal_lj_expand_ext.o: $(ALL_H) lal_lj_expand.h lal_lj_expand_ext.cpp lal_base_atomic.h

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



$(OBJ_DIR)/cg_cmm.cubin: lal_cg_cmm.cu lal_precision.h lal_preprocessor.h

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

$(OBJ_DIR)/lj_sdk.cubin: lal_lj_sdk.cu lal_precision.h lal_preprocessor.h

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



$(OBJ_DIR)/cg_cmm_cubin.h: $(OBJ_DIR)/cg_cmm.cubin $(OBJ_DIR)/cg_cmm.cubin

	$(BIN2C) -c -n cg_cmm $(OBJ_DIR)/cg_cmm.cubin > $(OBJ_DIR)/cg_cmm_cubin.h

$(OBJ_DIR)/lj_sdk_cubin.h: $(OBJ_DIR)/lj_sdk.cubin $(OBJ_DIR)/lj_sdk.cubin

	$(BIN2C) -c -n lj_sdk $(OBJ_DIR)/lj_sdk.cubin > $(OBJ_DIR)/lj_sdk_cubin.h



$(OBJ_DIR)/lal_cg_cmm.o: $(ALL_H) lal_cg_cmm.h lal_cg_cmm.cpp $(OBJ_DIR)/cg_cmm_cubin.h $(OBJ_DIR)/lal_base_atomic.o

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

$(OBJ_DIR)/lal_lj_sdk.o: $(ALL_H) lal_lj_sdk.h lal_lj_sdk.cpp $(OBJ_DIR)/lj_sdk_cubin.h $(OBJ_DIR)/lal_base_atomic.o

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



$(OBJ_DIR)/lal_cg_cmm_ext.o: $(ALL_H) lal_cg_cmm.h lal_cg_cmm_ext.cpp lal_base_atomic.h

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

$(OBJ_DIR)/lal_lj_sdk_ext.o: $(ALL_H) lal_lj_sdk.h lal_lj_sdk_ext.cpp lal_base_atomic.h

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



$(OBJ_DIR)/cg_cmm_long.cubin: lal_cg_cmm_long.cu lal_precision.h lal_preprocessor.h

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

$(OBJ_DIR)/lj_sdk_long.cubin: lal_lj_sdk_long.cu lal_precision.h lal_preprocessor.h

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



$(OBJ_DIR)/cg_cmm_long_cubin.h: $(OBJ_DIR)/cg_cmm_long.cubin $(OBJ_DIR)/cg_cmm_long.cubin

	$(BIN2C) -c -n cg_cmm_long $(OBJ_DIR)/cg_cmm_long.cubin > $(OBJ_DIR)/cg_cmm_long_cubin.h

$(OBJ_DIR)/lj_sdk_long_cubin.h: $(OBJ_DIR)/lj_sdk_long.cubin $(OBJ_DIR)/lj_sdk_long.cubin

	$(BIN2C) -c -n lj_sdk_long $(OBJ_DIR)/lj_sdk_long.cubin > $(OBJ_DIR)/lj_sdk_long_cubin.h



$(OBJ_DIR)/lal_cg_cmm_long.o: $(ALL_H) lal_cg_cmm_long.h lal_cg_cmm_long.cpp $(OBJ_DIR)/cg_cmm_long_cubin.h $(OBJ_DIR)/lal_base_atomic.o

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

$(OBJ_DIR)/lal_lj_sdk_long.o: $(ALL_H) lal_lj_sdk_long.h lal_lj_sdk_long.cpp $(OBJ_DIR)/lj_sdk_long_cubin.h $(OBJ_DIR)/lal_base_atomic.o

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



$(OBJ_DIR)/lal_cg_cmm_long_ext.o: $(ALL_H) lal_cg_cmm_long.h lal_cg_cmm_long_ext.cpp lal_base_charge.h

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

$(OBJ_DIR)/lal_lj_sdk_long_ext.o: $(ALL_H) lal_lj_sdk_long.h lal_lj_sdk_long_ext.cpp lal_base_charge.h

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



$(OBJ_DIR)/eam.cubin: lal_eam.cu lal_precision.h lal_preprocessor.h

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

lib/gpu/Opencl.makefile

+16 −16
Original line number Diff line number Diff line
@@ -32,8 +32,8 @@ OBJS = $(OBJ_DIR)/lal_atom.o $(OBJ_DIR)/lal_answer.o \ 
       $(OBJ_DIR)/lal_coul_long.o $(OBJ_DIR)/lal_coul_long_ext.o \

       $(OBJ_DIR)/lal_morse.o $(OBJ_DIR)/lal_morse_ext.o \

       $(OBJ_DIR)/lal_charmm_long.o $(OBJ_DIR)/lal_charmm_long_ext.o \

       $(OBJ_DIR)/lal_cg_cmm.o $(OBJ_DIR)/lal_cg_cmm_ext.o \

       $(OBJ_DIR)/lal_cg_cmm_long.o $(OBJ_DIR)/lal_cg_cmm_long_ext.o \

       $(OBJ_DIR)/lal_lj_sdk.o $(OBJ_DIR)/lal_lj_sdk_ext.o \

       $(OBJ_DIR)/lal_lj_sdk_long.o $(OBJ_DIR)/lal_lj_sdk_long_ext.o \

       $(OBJ_DIR)/lal_eam.o $(OBJ_DIR)/lal_eam_ext.o \

       $(OBJ_DIR)/lal_eam_fs_ext.o $(OBJ_DIR)/lal_eam_alloy_ext.o \

       $(OBJ_DIR)/lal_buck.o $(OBJ_DIR)/lal_buck_ext.o \
@@ -75,8 +75,8 @@ KERS = $(OBJ_DIR)/device_cl.h $(OBJ_DIR)/atom_cl.h \ 
       $(OBJ_DIR)/lj_coul_long_cl.h $(OBJ_DIR)/lj_dsf_cl.h \

       $(OBJ_DIR)/lj_class2_long_cl.h \

       $(OBJ_DIR)/coul_long_cl.h $(OBJ_DIR)/morse_cl.h \

       $(OBJ_DIR)/charmm_long_cl.h $(OBJ_DIR)/cg_cmm_cl.h \

       $(OBJ_DIR)/cg_cmm_long_cl.h $(OBJ_DIR)/neighbor_gpu_cl.h \

       $(OBJ_DIR)/charmm_long_cl.h $(OBJ_DIR)/lj_sdk_cl.h \

       $(OBJ_DIR)/lj_sdk_long_cl.h $(OBJ_DIR)/neighbor_gpu_cl.h \

       $(OBJ_DIR)/eam_cl.h $(OBJ_DIR)/buck_cl.h \

       $(OBJ_DIR)/buck_coul_cl.h $(OBJ_DIR)/buck_coul_long_cl.h \

       $(OBJ_DIR)/table_cl.h $(OBJ_DIR)/yukawa_cl.h \
@@ -273,23 +273,23 @@ $(OBJ_DIR)/lal_lj_expand.o: $(ALL_H) lal_lj_expand.h lal_lj_expand.cpp $(OBJ_DI 
$(OBJ_DIR)/lal_lj_expand_ext.o: $(ALL_H) lal_lj_expand.h lal_lj_expand_ext.cpp lal_base_atomic.h

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



$(OBJ_DIR)/cg_cmm_cl.h: lal_cg_cmm.cu $(PRE1_H)

	$(BSH) ./geryon/file_to_cstr.sh cg_cmm $(PRE1_H) lal_cg_cmm.cu $(OBJ_DIR)/cg_cmm_cl.h;

$(OBJ_DIR)/lj_sdk_cl.h: lal_lj_sdk.cu $(PRE1_H)

	$(BSH) ./geryon/file_to_cstr.sh lj_sdk $(PRE1_H) lal_lj_sdk.cu $(OBJ_DIR)/lj_sdk_cl.h;



$(OBJ_DIR)/lal_cg_cmm.o: $(ALL_H) lal_cg_cmm.h lal_cg_cmm.cpp  $(OBJ_DIR)/cg_cmm_cl.h $(OBJ_DIR)/cg_cmm_cl.h $(OBJ_DIR)/lal_base_atomic.o

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

$(OBJ_DIR)/lal_lj_sdk.o: $(ALL_H) lal_lj_sdk.h lal_lj_sdk.cpp  $(OBJ_DIR)/lj_sdk_cl.h $(OBJ_DIR)/lj_sdk_cl.h $(OBJ_DIR)/lal_base_atomic.o

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



$(OBJ_DIR)/lal_cg_cmm_ext.o: $(ALL_H) lal_cg_cmm.h lal_cg_cmm_ext.cpp lal_base_atomic.h

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

$(OBJ_DIR)/lal_lj_sdk_ext.o: $(ALL_H) lal_lj_sdk.h lal_lj_sdk_ext.cpp lal_base_atomic.h

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



$(OBJ_DIR)/cg_cmm_long_cl.h: lal_cg_cmm_long.cu $(PRE1_H)

	$(BSH) ./geryon/file_to_cstr.sh cg_cmm_long $(PRE1_H) lal_cg_cmm_long.cu $(OBJ_DIR)/cg_cmm_long_cl.h;

$(OBJ_DIR)/lj_sdk_long_cl.h: lal_lj_sdk_long.cu $(PRE1_H)

	$(BSH) ./geryon/file_to_cstr.sh lj_sdk_long $(PRE1_H) lal_lj_sdk_long.cu $(OBJ_DIR)/lj_sdk_long_cl.h;



$(OBJ_DIR)/lal_cg_cmm_long.o: $(ALL_H) lal_cg_cmm_long.h lal_cg_cmm_long.cpp  $(OBJ_DIR)/cg_cmm_long_cl.h $(OBJ_DIR)/cg_cmm_long_cl.h $(OBJ_DIR)/lal_base_atomic.o

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

$(OBJ_DIR)/lal_lj_sdk_long.o: $(ALL_H) lal_lj_sdk_long.h lal_lj_sdk_long.cpp  $(OBJ_DIR)/lj_sdk_long_cl.h $(OBJ_DIR)/lj_sdk_long_cl.h $(OBJ_DIR)/lal_base_atomic.o

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



$(OBJ_DIR)/lal_cg_cmm_long_ext.o: $(ALL_H) lal_cg_cmm_long.h lal_cg_cmm_long_ext.cpp lal_base_charge.h

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

$(OBJ_DIR)/lal_lj_sdk_long_ext.o: $(ALL_H) lal_lj_sdk_long.h lal_lj_sdk_long_ext.cpp lal_base_charge.h

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



$(OBJ_DIR)/eam_cl.h: lal_eam.cu $(PRE1_H)

	$(BSH) ./geryon/file_to_cstr.sh eam $(PRE1_H) lal_eam.cu $(OBJ_DIR)/eam_cl.h;

lib/gpu/lal_cg_cmm.cpp

+17 −17
Original line number Diff line number Diff line 
/***************************************************************************

                                 cg_cmm.cpp

                                 lj_sdk.cpp

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

                            W. Michael Brown (ORNL)
@@ -14,14 +14,14 @@ 
 ***************************************************************************/



#if defined(USE_OPENCL)

#include "cg_cmm_cl.h"

#include "lj_sdk_cl.h"

#elif defined(USE_CUDART)

const char *cg_cmm=0;

const char *lj_sdk=0;

#else

#include "cg_cmm_cubin.h"

#include "lj_sdk_cubin.h"

#endif



#include "lal_cg_cmm.h"

#include "lal_lj_sdk.h"

#include <cassert>

using namespace LAMMPS_AL;

#define CGCMMT CGCMM<numtyp, acctyp>
@@ -53,33 +53,33 @@ int CGCMMT::init(const int ntypes, double **host_cutsq, 
                          const double gpu_split, FILE *_screen) {

  int success;

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

                            _screen,cg_cmm,"k_cg_cmm");

                            _screen,lj_sdk,"k_lj_sdk");

  if (success!=0)

    return success;



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

  int cmm_types=ntypes;

  int sdk_types=ntypes;

  shared_types=false;

  int max_shared_types=this->device->max_shared_types();

  if (cmm_types<=max_shared_types && this->_block_size>=max_shared_types) {

    cmm_types=max_shared_types;

  if (sdk_types<=max_shared_types && this->_block_size>=max_shared_types) {

    sdk_types=max_shared_types;

    shared_types=true;

  }

  _cmm_types=cmm_types;

  _sdk_types=sdk_types;



  // Allocate a host write buffer for data initialization

  UCL_H_Vec<numtyp> host_write(cmm_types*cmm_types*32,*(this->ucl_device),

  UCL_H_Vec<numtyp> host_write(sdk_types*sdk_types*32,*(this->ucl_device),

                               UCL_WRITE_ONLY);



  for (int i=0; i<cmm_types*cmm_types; i++)

  for (int i=0; i<sdk_types*sdk_types; i++)

    host_write[i]=0.0;



  lj1.alloc(cmm_types*cmm_types,*(this->ucl_device),UCL_READ_ONLY);

  this->atom->type_pack4(ntypes,cmm_types,lj1,host_write,host_cutsq,

  lj1.alloc(sdk_types*sdk_types,*(this->ucl_device),UCL_READ_ONLY);

  this->atom->type_pack4(ntypes,sdk_types,lj1,host_write,host_cutsq,

                         host_cg_type,host_lj1,host_lj2);



  lj3.alloc(cmm_types*cmm_types,*(this->ucl_device),UCL_READ_ONLY);

  this->atom->type_pack4(ntypes,cmm_types,lj3,host_write,host_lj3,host_lj4,

  lj3.alloc(sdk_types*sdk_types,*(this->ucl_device),UCL_READ_ONLY);

  this->atom->type_pack4(ntypes,sdk_types,lj3,host_write,host_lj3,host_lj4,

                         host_offset);



  UCL_H_Vec<double> dview;
@@ -143,7 +143,7 @@ void CGCMMT::loop(const bool _eflag, const bool _vflag) { 
  } else {

    this->k_pair.set_size(GX,BX);

    this->k_pair.run(&this->atom->x, &lj1, &lj3,

                     &_cmm_types, &sp_lj, &this->nbor->dev_nbor,

                     &_sdk_types, &sp_lj, &this->nbor->dev_nbor,

                     &this->_nbor_data->begin(), &this->ans->force,

                     &this->ans->engv, &eflag, &vflag, &ainum,

                     &nbor_pitch, &this->_threads_per_atom);

lib/gpu/lal_cg_cmm.cu

+3 −3
Original line number Diff line number Diff line 
// **************************************************************************

//                                  cg_cmm.cu

//                                  lj_sdk.cu

//                             -------------------

//                           W. Michael Brown (ORNL)

//
@@ -24,7 +24,7 @@ texture<int4,1> pos_tex; 
#define pos_tex x_

#endif



__kernel void k_cg_cmm(const __global numtyp4 *restrict x_,

__kernel void k_lj_sdk(const __global numtyp4 *restrict x_,

                       const __global numtyp4 *restrict lj1,

                       const __global numtyp4 *restrict lj3,

                       const int lj_types,
@@ -116,7 +116,7 @@ __kernel void k_cg_cmm(const __global numtyp4 *restrict x_, 
  } // if ii

}



__kernel void k_cg_cmm_fast(const __global numtyp4 *restrict x_,

__kernel void k_lj_sdk_fast(const __global numtyp4 *restrict x_,

                            const __global numtyp4 *restrict lj1_in,

                            const __global numtyp4 *restrict lj3_in,

                            const __global numtyp *restrict sp_lj_in,

lib/gpu/lal_cg_cmm.h

+2 −2
Original line number Diff line number Diff line 
/***************************************************************************

                                  cg_cmm.h

                                  lj_sdk.h

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

                            W. Michael Brown (ORNL)
@@ -67,7 +67,7 @@ class CGCMM : public BaseAtomic<numtyp, acctyp> { 
  bool shared_types;



  /// Number of atom types

  int _cmm_types;

  int _sdk_types;



 private:

  bool _allocated;

CRA Git | Maintained and supported by SUSTech CRA and CCSE