diff --git a/lib/cuda/pair_cg_cmm_coul_cut_cuda.cu b/lib/cuda/pair_cg_cmm_coul_cut_cuda.cu deleted file mode 100644 index bafb83705a9dfb9a253051032986d009f0793edc..0000000000000000000000000000000000000000 --- a/lib/cuda/pair_cg_cmm_coul_cut_cuda.cu +++ /dev/null @@ -1,81 +0,0 @@ -/* ---------------------------------------------------------------------- - 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. -------------------------------------------------------------------------- */ - -#include <stdio.h> - -#define _lj1 MY_AP(coeff1) -#define _lj2 MY_AP(coeff2) -#define _lj3 MY_AP(coeff3) -#define _lj4 MY_AP(coeff4) -#define _cg_type MY_AP(coeff5) - - -#include "pair_cg_cmm_coul_cut_cuda_cu.h" -#include <time.h> - - - - -void Cuda_PairCGCMMCoulCutCuda_Init(cuda_shared_data* sdata) -{ - Cuda_Pair_Init_AllStyles(sdata, 5, true, false); - -} - - - - -void Cuda_PairCGCMMCoulCutCuda(cuda_shared_data* sdata, cuda_shared_neighlist* sneighlist, int eflag, int vflag, int eflag_atom, int vflag_atom) -{ - - // initialize only on first call - static short init = 0; - - if(! init) { - init = 1; - Cuda_PairCGCMMCoulCutCuda_Init(sdata); - } - - dim3 grid, threads; - int sharedperproc; - - Cuda_Pair_PreKernel_AllStyles(sdata, sneighlist, eflag, vflag, grid, threads, sharedperproc, true, 128); - - cudaStream_t* streams = (cudaStream_t*) CudaWrapper_returnStreams(); - - if(sdata->pair.use_block_per_atom) - Pair_Kernel_BpA<PAIR_CG_CMM, COUL_CUT, DATA_NONE> - <<< grid, threads, sharedperproc* sizeof(ENERGY_CFLOAT)*threads.x, streams[1]>>> (eflag, vflag, eflag_atom, vflag_atom); - else - Pair_Kernel_TpA<PAIR_CG_CMM, COUL_CUT, DATA_NONE> - <<< grid, threads, sharedperproc* sizeof(ENERGY_CFLOAT)*threads.x, streams[1]>>> (eflag, vflag, eflag_atom, vflag_atom); - - Cuda_Pair_PostKernel_AllStyles(sdata, grid, sharedperproc, eflag, vflag); -} - -#undef _lj1 -#undef _lj2 -#undef _lj3 -#undef _lj4 -#undef _cg_type - diff --git a/lib/cuda/pair_cg_cmm_coul_cut_cuda_cu.h b/lib/cuda/pair_cg_cmm_coul_cut_cuda_cu.h deleted file mode 100644 index fff7db7a646a72d9676983f27326a4f907d3cd5f..0000000000000000000000000000000000000000 --- a/lib/cuda/pair_cg_cmm_coul_cut_cuda_cu.h +++ /dev/null @@ -1,26 +0,0 @@ -/* ---------------------------------------------------------------------- - 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. -------------------------------------------------------------------------- */ - -#include "cuda_shared.h" - -extern "C" void Cuda_PairCGCMMCoulCutCuda(cuda_shared_data* sdata, cuda_shared_neighlist* sneighlist, int eflag, int vflag, int eflag_atom, int vflag_atom); diff --git a/lib/cuda/pair_cg_cmm_coul_debye_cuda.cu b/lib/cuda/pair_cg_cmm_coul_debye_cuda.cu deleted file mode 100644 index ee18b818990149da8b235b53c813a68a3ee2af3f..0000000000000000000000000000000000000000 --- a/lib/cuda/pair_cg_cmm_coul_debye_cuda.cu +++ /dev/null @@ -1,81 +0,0 @@ -/* ---------------------------------------------------------------------- - 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. -------------------------------------------------------------------------- */ - -#include <stdio.h> - -#define _lj1 MY_AP(coeff1) -#define _lj2 MY_AP(coeff2) -#define _lj3 MY_AP(coeff3) -#define _lj4 MY_AP(coeff4) -#define _cg_type MY_AP(coeff5) - - -#include "pair_cg_cmm_coul_debye_cuda_cu.h" -#include <time.h> - - - - -void Cuda_PairCGCMMCoulDebyeCuda_Init(cuda_shared_data* sdata) -{ - Cuda_Pair_Init_AllStyles(sdata, 5, true, false); - -} - - - - -void Cuda_PairCGCMMCoulDebyeCuda(cuda_shared_data* sdata, cuda_shared_neighlist* sneighlist, int eflag, int vflag, int eflag_atom, int vflag_atom) -{ - - // initialize only on first call - static short init = 0; - - if(! init) { - init = 1; - Cuda_PairCGCMMCoulDebyeCuda_Init(sdata); - } - - dim3 grid, threads; - int sharedperproc; - - Cuda_Pair_PreKernel_AllStyles(sdata, sneighlist, eflag, vflag, grid, threads, sharedperproc, true, 128); - - cudaStream_t* streams = (cudaStream_t*) CudaWrapper_returnStreams(); - - if(sdata->pair.use_block_per_atom) - Pair_Kernel_BpA<PAIR_CG_CMM, COUL_DEBYE, DATA_NONE> - <<< grid, threads, sharedperproc* sizeof(ENERGY_CFLOAT)*threads.x, streams[1]>>> (eflag, vflag, eflag_atom, vflag_atom); - else - Pair_Kernel_TpA<PAIR_CG_CMM, COUL_DEBYE, DATA_NONE> - <<< grid, threads, sharedperproc* sizeof(ENERGY_CFLOAT)*threads.x, streams[1]>>> (eflag, vflag, eflag_atom, vflag_atom); - - Cuda_Pair_PostKernel_AllStyles(sdata, grid, sharedperproc, eflag, vflag); -} - -#undef _lj1 -#undef _lj2 -#undef _lj3 -#undef _lj4 -#undef _cg_type - diff --git a/lib/cuda/pair_cg_cmm_coul_debye_cuda_cu.h b/lib/cuda/pair_cg_cmm_coul_debye_cuda_cu.h deleted file mode 100644 index 5b0b6af597f064ed9644f2920432540df9c1ecb7..0000000000000000000000000000000000000000 --- a/lib/cuda/pair_cg_cmm_coul_debye_cuda_cu.h +++ /dev/null @@ -1,26 +0,0 @@ -/* ---------------------------------------------------------------------- - 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. -------------------------------------------------------------------------- */ - -#include "cuda_shared.h" - -extern "C" void Cuda_PairCGCMMCoulDebyeCuda(cuda_shared_data* sdata, cuda_shared_neighlist* sneighlist, int eflag, int vflag, int eflag_atom, int vflag_atom); diff --git a/lib/cuda/pair_cg_cmm_coul_long_cuda.cu b/lib/cuda/pair_cg_cmm_coul_long_cuda.cu deleted file mode 100644 index 67c6ebaa7f3a275612648f3c7d8ad00dbbf3a765..0000000000000000000000000000000000000000 --- a/lib/cuda/pair_cg_cmm_coul_long_cuda.cu +++ /dev/null @@ -1,81 +0,0 @@ -/* ---------------------------------------------------------------------- - 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. -------------------------------------------------------------------------- */ - -#include <stdio.h> - -#define _lj1 MY_AP(coeff1) -#define _lj2 MY_AP(coeff2) -#define _lj3 MY_AP(coeff3) -#define _lj4 MY_AP(coeff4) -#define _cg_type MY_AP(coeff5) - - -#include "pair_cg_cmm_coul_long_cuda_cu.h" -#include <time.h> - - - - -void Cuda_PairCGCMMCoulLongCuda_Init(cuda_shared_data* sdata) -{ - Cuda_Pair_Init_AllStyles(sdata, 5, true, false); - -} - - - - -void Cuda_PairCGCMMCoulLongCuda(cuda_shared_data* sdata, cuda_shared_neighlist* sneighlist, int eflag, int vflag, int eflag_atom, int vflag_atom) -{ - - // initialize only on first call - static short init = 0; - - if(! init) { - init = 1; - Cuda_PairCGCMMCoulLongCuda_Init(sdata); - } - - dim3 grid, threads; - int sharedperproc; - - Cuda_Pair_PreKernel_AllStyles(sdata, sneighlist, eflag, vflag, grid, threads, sharedperproc, true, 128); - - cudaStream_t* streams = (cudaStream_t*) CudaWrapper_returnStreams(); - - if(sdata->pair.use_block_per_atom) - Pair_Kernel_BpA<PAIR_CG_CMM, COUL_LONG, DATA_NONE> - <<< grid, threads, sharedperproc* sizeof(ENERGY_CFLOAT)*threads.x, streams[1]>>> (eflag, vflag, eflag_atom, vflag_atom); - else - Pair_Kernel_TpA<PAIR_CG_CMM, COUL_LONG, DATA_NONE> - <<< grid, threads, sharedperproc* sizeof(ENERGY_CFLOAT)*threads.x, streams[1]>>> (eflag, vflag, eflag_atom, vflag_atom); - - Cuda_Pair_PostKernel_AllStyles(sdata, grid, sharedperproc, eflag, vflag); -} - -#undef _lj1 -#undef _lj2 -#undef _lj3 -#undef _lj4 -#undef _cg_type - diff --git a/lib/cuda/pair_cg_cmm_coul_long_cuda_cu.h b/lib/cuda/pair_cg_cmm_coul_long_cuda_cu.h deleted file mode 100644 index 86befd78b8705dcbe46fa3bd30252f23c5b613bd..0000000000000000000000000000000000000000 --- a/lib/cuda/pair_cg_cmm_coul_long_cuda_cu.h +++ /dev/null @@ -1,26 +0,0 @@ -/* ---------------------------------------------------------------------- - 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. -------------------------------------------------------------------------- */ - -#include "cuda_shared.h" - -extern "C" void Cuda_PairCGCMMCoulLongCuda(cuda_shared_data* sdata, cuda_shared_neighlist* sneighlist, int eflag, int vflag, int eflag_atom, int vflag_atom); diff --git a/lib/cuda/pair_cg_cmm_cuda.cu b/lib/cuda/pair_cg_cmm_cuda.cu deleted file mode 100644 index ccc4782ec04dd098cef42ced5de3d6196566bdf8..0000000000000000000000000000000000000000 --- a/lib/cuda/pair_cg_cmm_cuda.cu +++ /dev/null @@ -1,87 +0,0 @@ -/* ---------------------------------------------------------------------- - 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. -------------------------------------------------------------------------- */ - -#include <stdio.h> - -#define _lj1 MY_AP(coeff1) -#define _lj2 MY_AP(coeff2) -#define _lj3 MY_AP(coeff3) -#define _lj4 MY_AP(coeff4) -#define _cg_type MY_AP(coeff5) - -enum {CG_NOT_SET = 0, CG_LJ9_6, CG_LJ12_4, CG_LJ12_6, NUM_CG_TYPES, - CG_COUL_NONE, CG_COUL_CUT, CG_COUL_DEBYE, CG_COUL_LONG - }; - -#include "pair_cg_cmm_cuda_cu.h" -#include "pair_cg_cmm_cuda_kernel_nc.cu" -#include <time.h> - - - - -void Cuda_PairCGCMMCuda_Init(cuda_shared_data* sdata) -{ - Cuda_Pair_Init_AllStyles(sdata, 5, false, false); - -} - - - - -void Cuda_PairCGCMMCuda(cuda_shared_data* sdata, cuda_shared_neighlist* sneighlist, int eflag, int vflag, int eflag_atom, int vflag_atom) -{ - - // initialize only on first call - static short init = 0; - - if(! init) { - init = 1; - Cuda_PairCGCMMCuda_Init(sdata); - } - - dim3 grid, threads; - int sharedperproc; - - int maxthreads = 128; - - Cuda_Pair_PreKernel_AllStyles(sdata, sneighlist, eflag, vflag, grid, threads, sharedperproc, false, maxthreads); - - cudaStream_t* streams = (cudaStream_t*) CudaWrapper_returnStreams(); - - if(sdata->pair.use_block_per_atom) - Pair_Kernel_BpA<PAIR_CG_CMM, COUL_NONE, DATA_NONE> - <<< grid, threads, sharedperproc* sizeof(ENERGY_CFLOAT)*threads.x, streams[1]>>> (eflag, vflag, eflag_atom, vflag_atom); - else - Pair_Kernel_TpA<PAIR_CG_CMM, COUL_NONE, DATA_NONE> - <<< grid, threads, sharedperproc* sizeof(ENERGY_CFLOAT)*threads.x, streams[1]>>> (eflag, vflag, eflag_atom, vflag_atom); - - Cuda_Pair_PostKernel_AllStyles(sdata, grid, sharedperproc, eflag, vflag); -} - -#undef _lj1 -#undef _lj2 -#undef _lj3 -#undef _lj4 -#undef _cg_type - diff --git a/lib/cuda/pair_cg_cmm_cuda_cu.h b/lib/cuda/pair_cg_cmm_cuda_cu.h deleted file mode 100644 index 739c0ae28f2885015ef50a1048ab7c8e7a90d29c..0000000000000000000000000000000000000000 --- a/lib/cuda/pair_cg_cmm_cuda_cu.h +++ /dev/null @@ -1,26 +0,0 @@ -/* ---------------------------------------------------------------------- - 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. -------------------------------------------------------------------------- */ - -#include "cuda_shared.h" - -extern "C" void Cuda_PairCGCMMCuda(cuda_shared_data* sdata, cuda_shared_neighlist* sneighlist, int eflag, int vflag, int eflag_atom, int vflag_atom); diff --git a/lib/cuda/pair_cg_cmm_cuda_kernel_nc.cu b/lib/cuda/pair_cg_cmm_cuda_kernel_nc.cu deleted file mode 100644 index a19903b41abbe0761542275f9f0a59cd046c6549..0000000000000000000000000000000000000000 --- a/lib/cuda/pair_cg_cmm_cuda_kernel_nc.cu +++ /dev/null @@ -1,49 +0,0 @@ -/* ---------------------------------------------------------------------- - 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. -------------------------------------------------------------------------- */ - -__device__ inline F_CFLOAT PairCGCMMCuda_Eval(const F_CFLOAT &rsq, const int ij_type, F_CFLOAT &factor_lj, int &eflag, ENERGY_CFLOAT &evdwl) //0.11 of 0.4 -{ - const F_CFLOAT r2inv = F_F(1.0) / rsq; - const int cg_type = _cg_type[ij_type]; - const F_CFLOAT r4inv = r2inv * r2inv; - const F_CFLOAT rNinv_first = cg_type != CG_LJ9_6 ? r4inv : _RSQRT_(rsq); - const F_CFLOAT rNinv_second = cg_type != CG_LJ12_4 ? -r2inv : -F_F(1.0); - const F_CFLOAT forcelj = r4inv * (_lj1[ij_type] * r4inv * rNinv_first + _lj2[ij_type] * rNinv_second); - - if(eflag) evdwl += factor_lj * (r4inv * (_lj3[ij_type] * r4inv * rNinv_first + _lj4[ij_type] * rNinv_second) - _offset[ij_type]); - - return factor_lj * forcelj * r2inv; -} - -/*__device__ inline F_CFLOAT PairCGCMMCuda_Eval(const F_CFLOAT& rsq,const int ij_type,F_CFLOAT& factor_lj,int& eflag, ENERGY_CFLOAT& evdwl) -{ - const int cg_type = tex1Dfetch(_coeff5_gm_tex,ij_type); - const F_CFLOAT r2inv = F_F(1.0)/rsq; - const F_CFLOAT r4inv = r2inv*r2inv; - const F_CFLOAT rNinv_first = cg_type!=CG_LJ9_6?r4inv:_RSQRT_(rsq); - const F_CFLOAT rNinv_second = cg_type!=CG_LJ12_4?r2inv:F_F(1.0); - const F_CFLOAT forcelj = r4inv * (tex1Dfetch(_coeff1_gm_tex,ij_type)*r4inv*rNinv_first - tex1Dfetch(_coeff2_gm_tex,ij_type)*rNinv_second); - - if(eflag) evdwl += factor_lj*(r4inv*(tex1Dfetch(_coeff3_gm_tex,ij_type)*r4inv*rNinv_first-tex1Dfetch(_coeff4_gm_tex,ij_type)*rNinv_second)); - return factor_lj*forcelj*r2inv; -}*/