From bae4e45978cd3e161862b66d231be34a39579d4b Mon Sep 17 00:00:00 2001
From: Trung Nguyen <ndactrung@gmail.com>
Date: Sun, 13 Feb 2022 09:17:05 -0600
Subject: [PATCH] Added a comment to the added check while looping over the
 neighbors of ghost atoms

---
 lib/gpu/lal_sw.cu | 12 +++++++++---
 1 file changed, 9 insertions(+), 3 deletions(-)

diff --git a/lib/gpu/lal_sw.cu b/lib/gpu/lal_sw.cu
index c977fe0ca4..5c13caaa43 100644
--- a/lib/gpu/lal_sw.cu
+++ b/lib/gpu/lal_sw.cu
@@ -666,8 +666,11 @@ __kernel void k_sw_three_end(const __global numtyp4 *restrict x_,
         numtyp delr2z = kx.z - jx.z;
         numtyp rsq2 = delr2x*delr2x + delr2y*delr2y + delr2z*delr2z;
 
-        if (!(rsq2>(numtyp)0 && rsq2<sw_cut_ik*sw_cut_ik)) continue;
-
+        // still need this check because j and k can be ghost atoms
+        //   and short neighbor lists are built for local atoms only
+        const numtyp cutsq_jk=sw_cut_ik*sw_cut_ik;
+        if (!(rsq2>(numtyp)0 && rsq2<cutsq_jk)) continue;
+        
         #ifndef ONETYPE
         const numtyp sw_sigma_gamma_ik=cut_sig_gamma[mtypek].y;
         const int mtypejik=jtype*ntypes*ntypes+itype+ktype;
@@ -785,7 +788,10 @@ __kernel void k_sw_three_end_vatom(const __global numtyp4 *restrict x_,
         numtyp delr2z = kx.z - jx.z;
         numtyp rsq2 = delr2x*delr2x + delr2y*delr2y + delr2z*delr2z;
 
-        if (!(rsq2>(numtyp)0 && rsq2<sw_cut_ik*sw_cut_ik)) continue;
+        // still need this check because j and k can be ghost atoms
+        //   and short neighbor lists are built for local atoms only
+        const numtyp cutsq_jk=sw_cut_ik*sw_cut_ik;
+        if (!(rsq2>(numtyp)0 && rsq2<cutsq_jk)) continue;
 
         #ifndef ONETYPE
         const numtyp sw_sigma_gamma_ik=cut_sig_gamma[mtypek].y;