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

Rename local buffers 'red_acc' in lal_tersoff*

Browse Source 
Rename __local red_acc in lal_tersoff
This commit is contained in:
Vsevak
2020-08-23 22:21:31 +03:00
parent 1879106cc3
commit 0842911cd8
3 changed files with 12 additions and 12 deletions
Download Patch File Download Diff File Expand all files Collapse all files

6
lib/gpu/lal_tersoff.cu
View File

@ -709,7 +709,7 @@ __kernel void k_tersoff_three_end(const __global numtyp4 *restrict x_,
for (int i=0; i<6; i++)
virial[i]=(acctyp)0;
__local int red_acc[BLOCK_PAIR];
__local int ijnum_shared[BLOCK_PAIR];
__syncthreads();
@ -789,14 +789,14 @@ __kernel void k_tersoff_three_end(const __global numtyp4 *restrict x_,
k &= NEIGHMASK;
if (k == i) {
ijnum = nbor_k;
red_acc[m] = ijnum;
ijnum_shared[m] = ijnum;
break;
}
}
numtyp r1 = ucl_sqrt(rsq1);
numtyp r1inv = ucl_rsqrt(rsq1);
if (ijnum < 0) ijnum = red_acc[m];
if (ijnum < 0) ijnum = ijnum_shared[m];
// idx to zetaij is shifted by n_stride relative to ijnum in dev_short_nbor
int idx = ijnum;

12
lib/gpu/lal_tersoff_mod.cu
View File

@ -719,7 +719,7 @@ __kernel void k_tersoff_mod_three_end(const __global numtyp4 *restrict x_,
for (int i=0; i<6; i++)
virial[i]=(acctyp)0;
__local int red_acc[BLOCK_PAIR];
__local int ijnum_shared[BLOCK_PAIR];
__syncthreads();
@ -799,14 +799,14 @@ __kernel void k_tersoff_mod_three_end(const __global numtyp4 *restrict x_,
k &= NEIGHMASK;
if (k == i) {
ijnum = nbor_k;
red_acc[m] = ijnum;
ijnum_shared[m] = ijnum;
break;
}
}
numtyp r1 = ucl_sqrt(rsq1);
numtyp r1inv = ucl_rsqrt(rsq1);
if (ijnum < 0) ijnum = red_acc[m];
if (ijnum < 0) ijnum = ijnum_shared[m];
// idx to zetaij is shifted by n_stride relative to ijnum in dev_short_nbor
int idx = ijnum;
@ -957,7 +957,7 @@ __kernel void k_tersoff_mod_three_end_vatom(const __global numtyp4 *restrict x_,
for (int i=0; i<6; i++)
virial[i]=(acctyp)0;
__local int red_acc[BLOCK_PAIR];
__local int ijnum_shared[BLOCK_PAIR];
__syncthreads();
@ -1037,14 +1037,14 @@ __kernel void k_tersoff_mod_three_end_vatom(const __global numtyp4 *restrict x_,
k &= NEIGHMASK;
if (k == i) {
ijnum = nbor_k;
red_acc[m] = ijnum;
ijnum_shared[m] = ijnum;
break;
}
}
numtyp r1 = ucl_sqrt(rsq1);
numtyp r1inv = ucl_rsqrt(rsq1);
if (ijnum < 0) ijnum = red_acc[m];
if (ijnum < 0) ijnum = ijnum_shared[m];
// idx to zetaij is shifted by n_stride relative to ijnum in dev_short_nbor
int idx = ijnum;

6
lib/gpu/lal_tersoff_zbl.cu
View File

@ -729,7 +729,7 @@ __kernel void k_tersoff_zbl_three_end(const __global numtyp4 *restrict x_,
for (int i=0; i<6; i++)
virial[i]=(acctyp)0;
__local int red_acc[BLOCK_PAIR];
__local int ijnum_shared[BLOCK_PAIR];
__syncthreads();
@ -809,14 +809,14 @@ __kernel void k_tersoff_zbl_three_end(const __global numtyp4 *restrict x_,
k &= NEIGHMASK;
if (k == i) {
ijnum = nbor_k;
red_acc[m] = ijnum;
ijnum_shared[m] = ijnum;
break;
}
}
numtyp r1 = ucl_sqrt(rsq1);
numtyp r1inv = ucl_rsqrt(rsq1);
if (ijnum < 0) ijnum = red_acc[m];
if (ijnum < 0) ijnum = ijnum_shared[m];
// idx to zetaij is shifted by n_stride relative to ijnum in dev_short_nbor
int idx = ijnum;