diff --git a/lib/gpu/lal_tersoff.cu b/lib/gpu/lal_tersoff.cu index 7d3ede8dfc..b08fddfd6e 100644 --- a/lib/gpu/lal_tersoff.cu +++ b/lib/gpu/lal_tersoff.cu @@ -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; diff --git a/lib/gpu/lal_tersoff_mod.cu b/lib/gpu/lal_tersoff_mod.cu index bb5aabca44..0f45653264 100644 --- a/lib/gpu/lal_tersoff_mod.cu +++ b/lib/gpu/lal_tersoff_mod.cu @@ -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; diff --git a/lib/gpu/lal_tersoff_zbl.cu b/lib/gpu/lal_tersoff_zbl.cu index cd8d453b3c..f631cab91f 100644 --- a/lib/gpu/lal_tersoff_zbl.cu +++ b/lib/gpu/lal_tersoff_zbl.cu @@ -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;