From e727ec1eace97dfb55bbe2a44c134a5bfbf0ecf1 Mon Sep 17 00:00:00 2001
From: Trung Nguyen <ndactrung@gmail.com>
Date: Sat, 9 Dec 2023 16:46:10 -0600
Subject: [PATCH] Fixed compiling bugs revealed by CUDA builds, removed unused
 variables

---
 lib/gpu/lal_sph_heatconduction.cu |  6 +-----
 lib/gpu/lal_sph_lj.cu             | 24 +++++++++---------------
 lib/gpu/lal_sph_taitwater.cu      | 10 ----------
 3 files changed, 10 insertions(+), 30 deletions(-)

diff --git a/lib/gpu/lal_sph_heatconduction.cu b/lib/gpu/lal_sph_heatconduction.cu
index 4cea1433e3..21c936347a 100644
--- a/lib/gpu/lal_sph_heatconduction.cu
+++ b/lib/gpu/lal_sph_heatconduction.cu
@@ -160,10 +160,6 @@ __kernel void k_sph_heatconduction_fast(const __global numtyp4 *restrict x_,
 
   #ifndef ONETYPE
   __local numtyp4 coeff[MAX_SHARED_TYPES*MAX_SHARED_TYPES];
-  __local numtyp sp_lj[4];
-  if (tid<4) {
-    sp_lj[tid]=sp_lj_in[tid];
-  }
   if (tid<MAX_SHARED_TYPES*MAX_SHARED_TYPES) {
     coeff[tid]=coeff_in[tid];
   }
@@ -207,7 +203,7 @@ __kernel void k_sph_heatconduction_fast(const __global numtyp4 *restrict x_,
       int jtype = jx.w;
       #ifndef ONETYPE
       int mtype=itype+jx.w;
-      const numtyp cutsq_p=cutsq[mtype];
+      const numtyp cutsq_p=coeff[mtype].z;
       #endif
 
       // Compute r12
diff --git a/lib/gpu/lal_sph_lj.cu b/lib/gpu/lal_sph_lj.cu
index b9fdc5e433..23863b5e28 100644
--- a/lib/gpu/lal_sph_lj.cu
+++ b/lib/gpu/lal_sph_lj.cu
@@ -59,7 +59,7 @@ _texture_2d( vel_tex,int4);
    return p = pc[0], c = pc[1]
 */
 
-void LJEOS2(const numtyp rho, const numtyp e, const numtyp cv, numtyp pc[2])
+ucl_inline void LJEOS2(const numtyp rho, const numtyp e, const numtyp cv, numtyp pc[2])
 {
   numtyp T = e/cv;
   numtyp beta = ucl_recip(T); // (numtyp)1.0 / T;
@@ -176,7 +176,7 @@ __kernel void k_sph_lj(const __global numtyp4 *restrict x_,
         numtyp cvj = extraj.z;
 
         numtyp h = coeffy; // cut[itype][jtype]
-        ih = ucl_recip(h); // (numtyp)1.0 / h;
+        numtyp ih = ucl_recip(h); // (numtyp)1.0 / h;
         numtyp ihsq = ih * ih;
         numtyp ihcub = ihsq * ih;
 
@@ -210,13 +210,13 @@ __kernel void k_sph_lj(const __global numtyp4 *restrict x_,
 
         // artificial viscosity (Monaghan 1992)
         numtyp fvisc = (numtyp)0;
-        if (delVdotDelR < (numyp)0) {
-          numtyp mu = h * delVdotDelR / (rsq + (numyp)0.01 * h * h);
+        if (delVdotDelR < (numtyp)0) {
+          numtyp mu = h * delVdotDelR / (rsq + (numtyp)0.01 * h * h);
           fvisc = -coeffx * (ci + cj) * mu / (rhoi + rhoj); // viscosity[itype][jtype]
         }
 
         // total pair force & thermal energy increment
-        numtyp force = -mass_itype * mass_type * (fi + fj + fvisc) * wfd;
+        numtyp force = -mass_itype * mass_jtype * (fi + fj + fvisc) * wfd;
         numtyp deltaE = (numtyp)-0.5 * force * delVdotDelR;
 
         f.x+=delx*force;
@@ -268,10 +268,6 @@ __kernel void k_sph_lj_fast(const __global numtyp4 *restrict x_,
 
   #ifndef ONETYPE
   __local numtyp4 coeff[MAX_SHARED_TYPES*MAX_SHARED_TYPES];
-  __local numtyp sp_lj[4];
-  if (tid<4) {
-    sp_lj[tid]=sp_lj_in[tid];
-  }
   if (tid<MAX_SHARED_TYPES*MAX_SHARED_TYPES) {
     coeff[tid]=coeff_in[tid];
   }
@@ -307,7 +303,6 @@ __kernel void k_sph_lj_fast(const __global numtyp4 *restrict x_,
     int itype=fast_mul((int)MAX_SHARED_TYPES,iw);
     #endif
     numtyp4 iv; fetch4(iv,i,vel_tex); //v_[i];
-    int itag=iv.w;
 
     const numtyp4 extrai = extra[i];
     numtyp rhoi = extrai.x;
@@ -333,10 +328,9 @@ __kernel void k_sph_lj_fast(const __global numtyp4 *restrict x_,
       int jtype = jx.w;
       #ifndef ONETYPE
       int mtype=itype+jx.w;
-      const numtyp cutsq_p=cutsq[mtype];
+      const numtyp cutsq_p=coeff[mtype].z; // cutsq[itype][jtype];
       #endif
       numtyp4 jv; fetch4(jv,j,vel_tex); //v_[j];
-      int jtag=jv.w;
 
       // Compute r12
       numtyp delx = ix.x-jx.x;
@@ -356,7 +350,7 @@ __kernel void k_sph_lj_fast(const __global numtyp4 *restrict x_,
         numtyp cvj = extraj.z;
 
         numtyp h = coeffy; // cut[itype][jtype]
-        ih = ih = ucl_recip(h); // (numtyp)1.0 / h;
+        numtyp ih = ucl_recip(h); // (numtyp)1.0 / h;
         numtyp ihsq = ih * ih;
         numtyp ihcub = ihsq * ih;
 
@@ -390,8 +384,8 @@ __kernel void k_sph_lj_fast(const __global numtyp4 *restrict x_,
 
         // artificial viscosity (Monaghan 1992)
         numtyp fvisc = (numtyp)0;
-        if (delVdotDelR < (numyp)0) {
-          numtyp mu = h * delVdotDelR / (rsq + (numyp)0.01 * h * h);
+        if (delVdotDelR < (numtyp)0) {
+          numtyp mu = h * delVdotDelR / (rsq + (numtyp)0.01 * h * h);
           fvisc = -coeffx * (ci + cj) * mu / (rhoi + rhoj); // viscosity[itype][jtype]
         }
 
diff --git a/lib/gpu/lal_sph_taitwater.cu b/lib/gpu/lal_sph_taitwater.cu
index cc32e4a3c9..708d3ae43b 100644
--- a/lib/gpu/lal_sph_taitwater.cu
+++ b/lib/gpu/lal_sph_taitwater.cu
@@ -136,7 +136,6 @@ __kernel void k_sph_taitwater(const __global numtyp4 *restrict x_,
         numtyp h = coeffy; // cut[itype][jtype]
         numtyp ih = ucl_recip(h); // (numtyp)1.0 / h;
         numtyp ihsq = ih * ih;
-        numtyp ihcub = ihsq * ih;
 
         numtyp wfd = h - ucl_sqrt(rsq);
         if (dimension == 3) {
@@ -222,10 +221,6 @@ __kernel void k_sph_taitwater_fast(const __global numtyp4 *restrict x_,
   #ifndef ONETYPE
   __local numtyp4 coeff[MAX_SHARED_TYPES*MAX_SHARED_TYPES];
   __local numtyp4 coeff2[MAX_SHARED_TYPES];
-  __local numtyp sp_lj[4];
-  if (tid<4) {
-    sp_lj[tid]=sp_lj_in[tid];
-  }
   if (tid<MAX_SHARED_TYPES) {
     coeff2[tid] = coeff2_in[tid];
   }
@@ -267,11 +262,9 @@ __kernel void k_sph_taitwater_fast(const __global numtyp4 *restrict x_,
     int itype=fast_mul((int)MAX_SHARED_TYPES,iw);
     #endif
     numtyp4 iv; fetch4(iv,i,vel_tex); //v_[i];
-    int itag=iv.w;
 
     const numtyp4 extrai = extra[i];
     numtyp rhoi = extrai.x;
-    numtyp massi= extrai.y;
 
     // compute pressure of atom i with Tait EOS
     numtyp tmp = rhoi / rho0_itype;
@@ -294,7 +287,6 @@ __kernel void k_sph_taitwater_fast(const __global numtyp4 *restrict x_,
       const numtyp cutsq_p=coeff[mtype].z;
       #endif
       numtyp4 jv; fetch4(jv,j,vel_tex); //v_[j];
-      int jtag=jv.w;
 
       // Compute r12
       numtyp delx = ix.x-jx.x;
@@ -315,12 +307,10 @@ __kernel void k_sph_taitwater_fast(const __global numtyp4 *restrict x_,
 
         const numtyp4 extraj = extra[j];
         numtyp rhoj = extraj.x;
-        numtyp massj= extraj.y;
 
         numtyp h = coeffy; // cut[itype][jtype]
         numtyp ih = ucl_recip(h); // (numtyp)1.0 / h;
         numtyp ihsq = ih * ih;
-        numtyp ihcub = ihsq * ih;
 
         numtyp wfd = h - ucl_sqrt(rsq);
         if (dimension == 3) {