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 #define MY_PREFIX fix_viscous_cuda
26 #include "cuda_shared.h"
27 #include "cuda_common.h"
28 #include "crm_cuda_utils.cu"
29
30 #include "fix_viscous_cuda_cu.h"
31 #include "fix_viscous_cuda_kernel.cu"
32
Cuda_FixViscousCuda_UpdateNmax(cuda_shared_data * sdata)33 void Cuda_FixViscousCuda_UpdateNmax(cuda_shared_data* sdata)
34 {
35 cudaMemcpyToSymbol(MY_AP(mask) , & sdata->atom.mask .dev_data, sizeof(int*));
36 cudaMemcpyToSymbol(MY_AP(nlocal) , & sdata->atom.nlocal , sizeof(int));
37 cudaMemcpyToSymbol(MY_AP(nmax) , & sdata->atom.nmax , sizeof(int));
38 cudaMemcpyToSymbol(MY_AP(v) , & sdata->atom.x .dev_data, sizeof(X_FLOAT*));
39 cudaMemcpyToSymbol(MY_AP(f) , & sdata->atom.f .dev_data, sizeof(F_FLOAT*));
40 cudaMemcpyToSymbol(MY_AP(type) , & sdata->atom.type .dev_data, sizeof(int*));
41 }
42
Cuda_FixViscousCuda_Init(cuda_shared_data * sdata)43 void Cuda_FixViscousCuda_Init(cuda_shared_data* sdata)
44 {
45 Cuda_FixViscousCuda_UpdateNmax(sdata);
46
47 }
48
49
Cuda_FixViscousCuda_PostForce(cuda_shared_data * sdata,int groupbit,void * gamma)50 void Cuda_FixViscousCuda_PostForce(cuda_shared_data* sdata, int groupbit, void* gamma)
51 {
52 if(sdata->atom.update_nmax)
53 Cuda_FixViscousCuda_UpdateNmax(sdata);
54
55 if(sdata->atom.update_nlocal)
56 cudaMemcpyToSymbol(MY_AP(nlocal) , & sdata->atom.nlocal , sizeof(int));
57
58
59 int3 layout = getgrid(sdata->atom.nlocal, 0);
60 dim3 threads(layout.z, 1, 1);
61 dim3 grid(layout.x, layout.y, 1);
62
63 Cuda_FixViscousCuda_PostForce_Kernel <<< grid, threads, 0>>> (groupbit, (F_FLOAT*) gamma);
64 cudaThreadSynchronize();
65 CUT_CHECK_ERROR("Cuda_Cuda_FixViscousCuda_PostForce: Kernel execution failed");
66
67 }
68