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;