Homec4science
Diffusion lammps (lammm-master) Browse

lammps/lib/cuda/fix_nve_cuda_kernel.cu
860bca7d5e5dlammm-master

Last Change
Raw File

fix_nve_cuda_kernel.cu
View Options

/* ----------------------------------------------------------------------
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.
------------------------------------------------------------------------- */
static inline __device__ void check_distance(X_FLOAT &xtmp,X_FLOAT &ytmp,X_FLOAT &ztmp,int &i,int groupbit)
{
if(_dist_check)
{
X_FLOAT tmp=xtmp-_xhold[i];
X_FLOAT d=tmp*tmp;
tmp=ytmp-_xhold[i+_maxhold];
d+=tmp*tmp;
tmp=ztmp-_xhold[i+2*_maxhold];
d+=tmp*tmp;
d=((i < _nlocal) && (_mask[i] & groupbit))?d:X_F(0.0);
if(not __all(d<=_triggerneighsq))
_reneigh_flag[0]=1;
}
}
__global__ void FixNVECuda_InitialIntegrate_Kernel(int groupbit)
{
X_FLOAT xtmp,ytmp,ztmp;
#ifdef CUDA_USE_BINNING
const unsigned bin = gridDim.y * blockIdx.x + blockIdx.y;
if(threadIdx.x < _bin_count_local[bin])
{
const int i = 3*blockDim.x * bin + threadIdx.x;
if(_mask[i] & groupbit)
{
F_FLOAT* my_f = _binned_f + i;
V_FLOAT* my_v = _binned_v + i;
X_FLOAT* my_x = _binned_x + i;
V_FLOAT dtfm = _dtf
if(_rmass_flag) dtfm*= V_F(1.0) / _binned_rmass[i];
else dtfm*= V_F(1.0) / _mass[_binned_type[blockDim.x * bin + threadIdx.x]];
V_FLOAT v_mem;
v_mem = *my_v += dtfm * (*my_f); xtmp = *my_x += _dtv * v_mem; my_f += blockDim.x; my_v += blockDim.x; my_x += blockDim.x;
v_mem = *my_v += dtfm * (*my_f); ytmp = *my_x += _dtv * v_mem; my_f += blockDim.x; my_v += blockDim.x; my_x += blockDim.x;
v_mem = *my_v += dtfm * (*my_f); ztmp = *my_x += _dtv * v_mem;
}
}
#else
int i=(blockIdx.x*gridDim.y+blockIdx.y)*blockDim.x+threadIdx.x;
if(i < _nlocal && _mask[i] & groupbit)
{
F_FLOAT* my_f = _f + i;
V_FLOAT* my_v = _v + i;
X_FLOAT* my_x = _x + i;
V_FLOAT dtfm = _dtf;
if(_rmass_flag) dtfm*= V_F(1.0) / _rmass[i];
else dtfm*= V_F(1.0) / _mass[_type[i]];
V_FLOAT v_mem;
v_mem = *my_v += dtfm * (*my_f); xtmp=*my_x += _dtv * v_mem; my_f += _nmax; my_v += _nmax; my_x += _nmax;
v_mem = *my_v += dtfm * (*my_f); ytmp=*my_x += _dtv * v_mem; my_f += _nmax; my_v += _nmax; my_x += _nmax;
v_mem = *my_v += dtfm * (*my_f); ztmp=*my_x += _dtv * v_mem;
}
#endif
check_distance(xtmp,ytmp,ztmp,i,groupbit);
}
__global__ void FixNVECuda_FinalIntegrate_Kernel(int groupbit)
{
#ifdef CUDA_USE_BINNING
const unsigned bin = gridDim.y * blockIdx.x + blockIdx.y;
if(threadIdx.x < _bin_count_local[bin])
{
const int i = 3*blockDim.x * bin + threadIdx.x;
if(_mask[i] & groupbit)
{
F_FLOAT* my_f = _binned_f + i;
V_FLOAT* my_v = _binned_v + i;
V_FLOAT dtfm = _dtf
if(_rmass_flag) dtfm*= V_F(1.0) / _binned_rmass[i];
else dtfm*= V_F(1.0) / _mass[_binned_type[blockDim.x * bin + threadIdx.x]];
*my_v += dtfm * (*my_f); my_f += blockDim.x; my_v += blockDim.x;
*my_v += dtfm * (*my_f); my_f += blockDim.x; my_v += blockDim.x;
*my_v += dtfm * (*my_f);
}
}
#else
int i=(blockIdx.x*gridDim.y+blockIdx.y)*blockDim.x+threadIdx.x;
if(i < _nlocal && _mask[i] & groupbit)
{
F_FLOAT* my_f = _f + i;
V_FLOAT* my_v = _v + i;
V_FLOAT dtfm = _dtf;
if(_rmass_flag) dtfm*=V_F(1.0) / _rmass[i];
else dtfm*=V_F(1.0) / _mass[_type[i]];
*my_v += dtfm * (*my_f); my_f += _nmax; my_v += _nmax;
*my_v += dtfm * (*my_f); my_f += _nmax; my_v += _nmax;
*my_v += dtfm * (*my_f);
}
#endif
}
c4science · Help