From ba49e7e400736005b915c7fad412fc6db64b7298 Mon Sep 17 00:00:00 2001
From: jpekkila <johannes.pekkila@aalto.fi>
Date: Mon, 7 Oct 2019 19:40:27 +0300
Subject: [PATCH] Replaced deprecated DCONST_INT calls with overloaded DCONST()

---
 src/core/device.cu              | 10 +++-------
 src/core/kernels/boundconds.cuh | 31 +++++++++++++++----------------
 2 files changed, 18 insertions(+), 23 deletions(-)

diff --git a/src/core/device.cu b/src/core/device.cu
index a6d409a..17e87f7 100644
--- a/src/core/device.cu
+++ b/src/core/device.cu
@@ -80,17 +80,13 @@ DCONST(const AcReal3Param param)
 {
     return d_mesh_info.real3_params[param];
 }
-constexpr VertexBufferHandle
+static __device__ constexpr VertexBufferHandle
 DCONST(const VertexBufferHandle handle)
 {
     return handle;
 }
-#define DCONST_INT(x) DCONST(x)
-#define DCONST_INT3(x) DCONST(x)
-#define DCONST_REAL(x) DCONST(x)
-#define DCONST_REAL3(x) DCONST(x)
-#define DEVICE_VTXBUF_IDX(i, j, k) ((i) + (j)*DCONST_INT(AC_mx) + (k)*DCONST_INT(AC_mxy))
-#define DEVICE_1D_COMPDOMAIN_IDX(i, j, k) ((i) + (j)*DCONST_INT(AC_nx) + (k)*DCONST_INT(AC_nxy))
+#define DEVICE_VTXBUF_IDX(i, j, k) ((i) + (j)*DCONST(AC_mx) + (k)*DCONST(AC_mxy))
+#define DEVICE_1D_COMPDOMAIN_IDX(i, j, k) ((i) + (j)*DCONST(AC_nx) + (k)*DCONST(AC_nxy))
 #define globalGridN (d_mesh_info.int3_params[AC_global_grid_n])
 //#define globalMeshM // Placeholder
 //#define localMeshN // Placeholder
diff --git a/src/core/kernels/boundconds.cuh b/src/core/kernels/boundconds.cuh
index 5c70114..1a01d1b 100644
--- a/src/core/kernels/boundconds.cuh
+++ b/src/core/kernels/boundconds.cuh
@@ -38,36 +38,35 @@ kernel_periodic_boundconds(const int3 start, const int3 end, AcReal* vtxbuf)
     if (i_dst >= end.x || j_dst >= end.y || k_dst >= end.z)
         return;
 
-    // if (i_dst >= DCONST_INT(AC_mx) || j_dst >= DCONST_INT(AC_my) || k_dst >= DCONST_INT(AC_mz))
+    // if (i_dst >= DCONST(AC_mx) || j_dst >= DCONST(AC_my) || k_dst >= DCONST(AC_mz))
     //    return;
 
     // If destination index is inside the computational domain, return since
     // the boundary conditions are only applied to the ghost zones
-    if (i_dst >= DCONST_INT(AC_nx_min) && i_dst < DCONST_INT(AC_nx_max) &&
-        j_dst >= DCONST_INT(AC_ny_min) && j_dst < DCONST_INT(AC_ny_max) &&
-        k_dst >= DCONST_INT(AC_nz_min) && k_dst < DCONST_INT(AC_nz_max))
+    if (i_dst >= DCONST(AC_nx_min) && i_dst < DCONST(AC_nx_max) && j_dst >= DCONST(AC_ny_min) &&
+        j_dst < DCONST(AC_ny_max) && k_dst >= DCONST(AC_nz_min) && k_dst < DCONST(AC_nz_max))
         return;
 
     // Find the source index
     // Map to nx, ny, nz coordinates
-    int i_src = i_dst - DCONST_INT(AC_nx_min);
-    int j_src = j_dst - DCONST_INT(AC_ny_min);
-    int k_src = k_dst - DCONST_INT(AC_nz_min);
+    int i_src = i_dst - DCONST(AC_nx_min);
+    int j_src = j_dst - DCONST(AC_ny_min);
+    int k_src = k_dst - DCONST(AC_nz_min);
 
     // Translate (s.t. the index is always positive)
-    i_src += DCONST_INT(AC_nx);
-    j_src += DCONST_INT(AC_ny);
-    k_src += DCONST_INT(AC_nz);
+    i_src += DCONST(AC_nx);
+    j_src += DCONST(AC_ny);
+    k_src += DCONST(AC_nz);
 
     // Wrap
-    i_src %= DCONST_INT(AC_nx);
-    j_src %= DCONST_INT(AC_ny);
-    k_src %= DCONST_INT(AC_nz);
+    i_src %= DCONST(AC_nx);
+    j_src %= DCONST(AC_ny);
+    k_src %= DCONST(AC_nz);
 
     // Map to mx, my, mz coordinates
-    i_src += DCONST_INT(AC_nx_min);
-    j_src += DCONST_INT(AC_ny_min);
-    k_src += DCONST_INT(AC_nz_min);
+    i_src += DCONST(AC_nx_min);
+    j_src += DCONST(AC_ny_min);
+    k_src += DCONST(AC_nz_min);
 
     const int src_idx = DEVICE_VTXBUF_IDX(i_src, j_src, k_src);
     const int dst_idx = DEVICE_VTXBUF_IDX(i_dst, j_dst, k_dst);