Skip to content
Snippets Groups Projects
Commit 20a71b1d authored by sjplimp's avatar sjplimp
Browse files 

git-svn-id: svn://svn.icms.temple.edu/lammps-ro/trunk@12640 f3b2605a-c512-4ea7-a41b-209d697bcdaa
parent 4ffd3808
No related branches found
No related tags found
No related merge requests found
Hide whitespace changes
Inline Side-by-side
Showing
with 0 additions and 483 deletions
lib/cuda/pair_cg_cmm_coul_cut_cuda.cu deleted 100644 → 0
+ 0
− 81
View file @ 4ffd3808
/* ----------------------------------------------------------------------
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
lib/cuda/pair_cg_cmm_coul_cut_cuda_cu.h deleted 100644 → 0
+ 0
− 26
View file @ 4ffd3808
/* ----------------------------------------------------------------------
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);
lib/cuda/pair_cg_cmm_coul_debye_cuda.cu deleted 100644 → 0
+ 0
− 81
View file @ 4ffd3808
/* ----------------------------------------------------------------------
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
lib/cuda/pair_cg_cmm_coul_debye_cuda_cu.h deleted 100644 → 0
+ 0
− 26
View file @ 4ffd3808
/* ----------------------------------------------------------------------
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);
lib/cuda/pair_cg_cmm_coul_long_cuda.cu deleted 100644 → 0
+ 0
− 81
View file @ 4ffd3808
/* ----------------------------------------------------------------------
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
lib/cuda/pair_cg_cmm_coul_long_cuda_cu.h deleted 100644 → 0
+ 0
− 26
View file @ 4ffd3808
/* ----------------------------------------------------------------------
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);
lib/cuda/pair_cg_cmm_cuda.cu deleted 100644 → 0
+ 0
− 87
View file @ 4ffd3808
/* ----------------------------------------------------------------------
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
lib/cuda/pair_cg_cmm_cuda_cu.h deleted 100644 → 0
+ 0
− 26
View file @ 4ffd3808
/* ----------------------------------------------------------------------
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);
lib/cuda/pair_cg_cmm_cuda_kernel_nc.cu deleted 100644 → 0
+ 0
− 49
View file @ 4ffd3808
/* ----------------------------------------------------------------------
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;
}*/
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment