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.
This commit is contained in:
@@ -40,14 +40,44 @@ typedef struct {
|
|||||||
} VertexBufferArray;
|
} VertexBufferArray;
|
||||||
|
|
||||||
__constant__ AcMeshInfo d_mesh_info;
|
__constant__ AcMeshInfo d_mesh_info;
|
||||||
#define DCONST_INT(X) (d_mesh_info.int_params[X])
|
static inline int __device__
|
||||||
#define DCONST_INT3(X) (d_mesh_info.int3_params[X])
|
DCONST(const AcIntParam param)
|
||||||
#define DCONST_REAL(X) (d_mesh_info.real_params[X])
|
{
|
||||||
#define DCONST_REAL3(X) (d_mesh_info.real3_params[X])
|
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_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_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 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_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/boundconds.cuh"
|
||||||
#include "kernels/integration.cuh"
|
#include "kernels/integration.cuh"
|
||||||
#include "kernels/reductions.cuh"
|
#include "kernels/reductions.cuh"
|
||||||
|
|||||||
Reference in New Issue
Block a user