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_gm)
27 #define _lj2 MY_AP(coeff2_gm)
28 #define _lj3 MY_AP(coeff3_gm)
29 #define _lj4 MY_AP(coeff4_gm)
30 #define _ljsw1 MY_AP(coeff5_gm)
31 #define _ljsw2 MY_AP(coeff6_gm)
32 #define _ljsw3 MY_AP(coeff7_gm)
33 #define _ljsw4 MY_AP(coeff8_gm)
34 #define _ljsw5 MY_AP(coeff9_gm)
35
36 #include "pair_lj_gromacs_cuda_cu.h"
37 #include "pair_lj_gromacs_cuda_kernel_nc.cu"
38
39 #include <time.h>
40
Cuda_PairLJGromacsCuda_Init(cuda_shared_data * sdata)41 void Cuda_PairLJGromacsCuda_Init(cuda_shared_data* sdata)
42 {
43 Cuda_Pair_Init_AllStyles(sdata, 9, false, true, true);
44 }
45
46
47
Cuda_PairLJGromacsCuda(cuda_shared_data * sdata,cuda_shared_neighlist * sneighlist,int eflag,int vflag,int eflag_atom,int vflag_atom)48 void Cuda_PairLJGromacsCuda(cuda_shared_data* sdata, cuda_shared_neighlist* sneighlist, int eflag, int vflag,
49 int eflag_atom, int vflag_atom)
50 {
51 static short init = 0;
52
53 if(! init) {
54 init = 1;
55 Cuda_PairLJGromacsCuda_Init(sdata);
56 }
57
58 dim3 grid, threads;
59 int sharedperproc;
60
61 Cuda_Pair_PreKernel_AllStyles(sdata, sneighlist, eflag, vflag, grid, threads, sharedperproc, false, 192);
62
63 cudaStream_t* streams = (cudaStream_t*) CudaWrapper_returnStreams();
64
65 if(sdata->pair.use_block_per_atom)
66 Pair_Kernel_BpA<PAIR_LJ_GROMACS, COUL_NONE, DATA_NONE>
67 <<< grid, threads, sharedperproc* sizeof(ENERGY_FLOAT)*threads.x, streams[1]>>> (eflag, vflag, eflag_atom, vflag_atom);
68 else
69 Pair_Kernel_TpA<PAIR_LJ_GROMACS, COUL_NONE, DATA_NONE>
70 <<< grid, threads, sharedperproc* sizeof(ENERGY_FLOAT)*threads.x, streams[1]>>> (eflag, vflag, eflag_atom, vflag_atom);
71
72 Cuda_Pair_PostKernel_AllStyles(sdata, grid, sharedperproc, eflag, vflag);
73
74 }
75
76 #undef _lj1
77 #undef _lj2
78 #undef _lj3
79 #undef _lj4
80 #undef _ljsw1
81 #undef _ljsw2
82 #undef _ljsw3
83 #undef _ljsw4
84 #undef _ljsw5
85