From d5b2e5bb42f42d405e1ab5ab8cd3454ee3316a32 Mon Sep 17 00:00:00 2001 From: jpekkila Date: Mon, 12 Aug 2019 14:05:35 +0300 Subject: [PATCH] Added placeholders for new built-in variables in the DSL. Also overloads to DCONST_INT etc. Naming still pending and old DCONST_REAL etc calls still work. --- src/core/device.cu | 38 ++++++++++++++++++++++++++++++++++---- 1 file changed, 34 insertions(+), 4 deletions(-) diff --git a/src/core/device.cu b/src/core/device.cu index 89bb6b1..9690278 100644 --- a/src/core/device.cu +++ b/src/core/device.cu @@ -40,14 +40,44 @@ typedef struct { } VertexBufferArray; __constant__ AcMeshInfo d_mesh_info; -#define DCONST_INT(X) (d_mesh_info.int_params[X]) -#define DCONST_INT3(X) (d_mesh_info.int3_params[X]) -#define DCONST_REAL(X) (d_mesh_info.real_params[X]) -#define DCONST_REAL3(X) (d_mesh_info.real3_params[X]) +static inline int __device__ +DCONST(const AcIntParam param) +{ + return d_mesh_info.int_params[param]; +} +static inline int3 __device__ +DCONST(const AcInt3Param param) +{ + return d_mesh_info.int3_params[param]; +} +static inline AcReal __device__ +DCONST(const AcRealParam param) +{ + return d_mesh_info.real_params[param]; +} +static inline AcReal3 __device__ +DCONST(const AcReal3Param param) +{ + return d_mesh_info.real3_params[param]; +} +#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 DCONST_INT(X) (d_mesh_info.int_params[X]) +//#define DCONST_INT3(X) (d_mesh_info.int3_params[X]) +//#define DCONST_REAL(X) (d_mesh_info.real_params[X]) +//#define DCONST_REAL3(X) (d_mesh_info.real3_params[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 globalGridN (d_mesh_info.int3_params[AC_global_grid_n]) +//#define globalMeshM // Placeholder +//#define localMeshN // Placeholder +//#define localMeshM // Placeholder +//#define localMeshN_min // Placeholder +//#define globalMeshN_min // Placeholder #define d_multigpu_offset (d_mesh_info.int3_params[AC_multigpu_offset]) +//#define d_multinode_offset (d_mesh_info.int3_params[AC_multinode_offset]) // Placeholder #include "kernels/boundconds.cuh" #include "kernels/integration.cuh" #include "kernels/reductions.cuh"