ICODE-ADC/lammps
4
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

git-svn-id: svn://svn.icms.temple.edu/lammps-ro/trunk@5906 f3b2605a-c512-4ea7-a41b-209d697bcdaa

Browse Source
This commit is contained in:
sjplimp
2011-04-07 16:07:43 +00:00
parent f3feb53326
commit 3480e40e5f
10 changed files with 105 additions and 130 deletions
Download Patch File Download Diff File Expand all files Collapse all files

28
lib/gpu/cmmc_long_gpu_kernel.cu
View File

@ -93,6 +93,10 @@ __inline float fetch_q(const int& i, const float *q)
#endif
#define SBBITS 30
#define NEIGHMASK 0x3FFFFFFF
__inline int sbmask(int j) { return j >> SBBITS & 3; }
__kernel void kernel_pair(__global numtyp4 *x_, __global numtyp4 *lj1,
__global numtyp4* lj3, const int lj_types,
__global numtyp *sp_lj_in, __global int *dev_nbor,
@ -139,14 +143,10 @@ __kernel void kernel_pair(__global numtyp4 *x_, __global numtyp4 *lj1,
int j=*nbor;
numtyp factor_lj, factor_coul;
if (j < nall) {
factor_lj = (numtyp)1.0;
factor_coul = (numtyp)0.0;
} else {
factor_lj = sp_lj[j/nall];
factor_coul = (numtyp)1.0-sp_lj[j/nall+4];
j %= nall;
}
factor_lj = sp_lj[sbmask(j)];
factor_coul = (numtyp)1.0-sp_lj[sbmask(j)+4];
j &= NEIGHMASK;
numtyp4 jx=fetch_pos(j,x_); //x_[j];
int jtype=jx.w;
@ -282,14 +282,10 @@ __kernel void kernel_pair_fast(__global numtyp4 *x_, __global numtyp4 *lj1_in,
int j=*nbor;
numtyp factor_lj, factor_coul;
if (j < nall) {
factor_lj = (numtyp)1.0;
factor_coul = (numtyp)0.0;
} else {
factor_lj = sp_lj[j/nall];
factor_coul = (numtyp)1.0-sp_lj[j/nall+4];
j %= nall;
}
factor_lj = sp_lj[sbmask(j)];
factor_coul = (numtyp)1.0-sp_lj[sbmask(j)+4];
j &= NEIGHMASK;
numtyp4 jx=fetch_pos(j,x_); //x_[j];
int mtype=itype+jx.w;