cwpearson/astaroth
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Replaced deprecated DCONST_INT calls with overloaded DCONST()

Browse Source
This commit is contained in:
jpekkila
2019-10-07 19:40:27 +03:00
parent 9c575f8059
commit ba49e7e400
2 changed files with 18 additions and 23 deletions
Download Patch File Download Diff File Expand all files Collapse all files

10
src/core/device.cu
View File

@@ -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

31
src/core/kernels/boundconds.cuh
View File

@@ -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);