1 /* ----------------------------------------------------------------------
2 LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator
3
4 Original Version:
5 http://lammps.sandia.gov, Sandia National Laboratories
6 Steve Plimpton, sjplimp@sandia.gov
7
8 See the README file in the top-level LAMMPS directory.
9
10 -----------------------------------------------------------------------
11
12 USER-CUDA Package and associated modifications:
13 https://sourceforge.net/projects/lammpscuda/
14
15 Christian Trott, christian.trott@tu-ilmenau.de
16 Lars Winterfeld, lars.winterfeld@tu-ilmenau.de
17 Theoretical Physics II, University of Technology Ilmenau, Germany
18
19 See the README file in the USER-CUDA directory.
20
21 This software is distributed under the GNU General Public License.
22 ------------------------------------------------------------------------- */
23
24 #include <stdio.h>
25
26 #define _lj1 MY_AP(coeff1)
27 #define _lj2 MY_AP(coeff2)
28 #define _lj3 MY_AP(coeff3)
29 #define _lj4 MY_AP(coeff4)
30
31 #include "pair_lj_class2_cuda_cu.h"
32 #include "pair_lj_class2_cuda_kernel_nc.cu"
33
34 #include <time.h>
35
Cuda_PairLJClass2Cuda_Init(cuda_shared_data * sdata)36 void Cuda_PairLJClass2Cuda_Init(cuda_shared_data* sdata)
37 {
38 Cuda_Pair_Init_AllStyles(sdata, 4);
39 }
40
Cuda_PairLJClass2Cuda(cuda_shared_data * sdata,cuda_shared_neighlist * sneighlist,int eflag,int vflag,int eflag_atom,int vflag_atom)41 void Cuda_PairLJClass2Cuda(cuda_shared_data* sdata, cuda_shared_neighlist* sneighlist, int eflag, int vflag, int eflag_atom, int vflag_atom)
42 {
43
44
45 static short init = 0;
46
47 if(! init) {
48 init = 1;
49 Cuda_PairLJClass2Cuda_Init(sdata);
50 }
51
52 dim3 grid, threads;
53 int sharedperproc;
54
55 //int maxthreads=192*sizeof(double)/sizeof(F_FLOAT);
56 //if(CUDA_ARCH==20) maxthreads*=2;
57 //cudaFuncSetCacheConfig(Pair_Kernel_TpA_opt<PAIR_LJ_CUT,COUL_NONE,DATA_NONE>,cudaFuncCachePreferL1);
58 Cuda_Pair_PreKernel_AllStyles(sdata, sneighlist, eflag, vflag, grid, threads, sharedperproc, false, 192);
59 cudaStream_t* streams = (cudaStream_t*) CudaWrapper_returnStreams();
60
61 if(sdata->pair.use_block_per_atom)
62 Pair_Kernel_BpA<PAIR_LJ_CLASS2, COUL_NONE, DATA_NONE>
63 <<< grid, threads, sharedperproc* sizeof(ENERGY_FLOAT)*threads.x, streams[1]>>> (eflag, vflag, eflag_atom, vflag_atom);
64 else
65 Pair_Kernel_TpA<PAIR_LJ_CLASS2, COUL_NONE, DATA_NONE>
66 <<< grid, threads, sharedperproc* sizeof(ENERGY_FLOAT)*threads.x, streams[1]>>> (eflag, vflag, eflag_atom, vflag_atom);
67
68 Cuda_Pair_PostKernel_AllStyles(sdata, grid, sharedperproc, eflag, vflag);
69 }
70
71
72 #undef _lj1
73 #undef _lj2
74 #undef _lj3
75 #undef _lj4
76