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

Periodic boundary conditions work with switchable system.

Browse Source 
Still some issue with the custom alternatives. Need to look into kernels again.
This commit is contained in:
Miikka Vaisala
2020-11-20 17:00:48 +08:00
parent 288693fab5
commit d4ee066b3c
3 changed files with 84 additions and 21 deletions
Download Patch File Download Diff File Expand all files Collapse all files

9
acc/mhd_solver/stencil_kernel.ac
View File

@@ -22,9 +22,12 @@ uniform int AC_save_steps;
uniform int AC_bin_steps; uniform int AC_bin_steps;
uniform int AC_start_step; uniform int AC_start_step;
// In3 params uniform int AC_bc_type_top_x;
uniform int3 AC_bc_type_bot; uniform int AC_bc_type_bot_x;
uniform int3 AC_bc_type_top; uniform int AC_bc_type_top_y;
uniform int AC_bc_type_bot_y;
uniform int AC_bc_type_top_z;
uniform int AC_bc_type_bot_z;
// Real params // Real params
uniform Scalar AC_dt; uniform Scalar AC_dt;

9
config/astaroth.conf
View File

@@ -13,6 +13,15 @@ AC_dsx = 0.04908738521
AC_dsy = 0.04908738521 AC_dsy = 0.04908738521
AC_dsz = 0.04908738521 AC_dsz = 0.04908738521
// 0 = periodic bc, 1 = symmetric bc, 2 = antisymmetric bc
AC_bc_type_top_x = 0
AC_bc_type_top_y = 0
AC_bc_type_top_z = 0
AC_bc_type_bot_x = 0
AC_bc_type_bot_y = 0
AC_bc_type_bot_z = 0
/* /*
* ============================================================================= * =============================================================================
* Run-time params * Run-time params

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

@@ -23,48 +23,84 @@ kernel_symmetric_boundconds(const int3 start, const int3 end, AcReal* vtxbuf, co
int i_src, j_src, k_src, boundloc; int i_src, j_src, k_src, boundloc;
int bsize = STENCIL_ORDER/(int) 2; int bsize = STENCIL_ORDER/(int) 2;
if (bindex.x != 0) //if (bindex.x != 0)
//{
// // Pick up the mirroring value.
// if ((i_dst < bsize) && ((bindex.x == 3) || (bindex.x ==1)))
// {
// boundloc = bsize; //Location of central border point.
// i_src = 2*boundloc - i_dst;
// } else if ((i_dst >= DCONST(AC_nx_min) - bsize) && ((bindex.x == 2) || (bindex.x ==1)))
// {
// boundloc = DCONST(AC_nx_min) - bsize - 1; //Location of central border point.
// i_src = 2*boundloc - i_dst;
// }
//}
//if (bindex.y != 0)
//{
// // Pick up the mirroring value.
// if ((j_dst < bsize) && ((bindex.y == 3) || (bindex.y ==1)))
// {
// boundloc = bsize; //Location of central border point.
// j_src = 2*boundloc - j_dst;
// } else if ((j_dst >= DCONST(AC_nx_min) - bsize) && ((bindex.y == 2) || (bindex.y ==1)))
// {
// boundloc = DCONST(AC_ny_min) - bsize - 1; //Location of central border point.
// i_src = 2*boundloc - j_dst;
// }
//}
//if (bindex.z != 0)
//{
// // Pick up the mirroring value.
// if ((k_dst < bsize) && ((bindex.z == 3) || (bindex.z ==1)))
// {
// boundloc = bsize; //Location of central border point.
// k_src = 2*boundloc - k_dst;
// } else if ((i_dst >= DCONST(AC_nz_min) - bsize) && ((bindex.z == 2) || (bindex.z ==1)))
// {
// boundloc = DCONST(AC_nz_min) - bsize - 1; //Location of central border point.
// k_src = 2*boundloc - k_dst;
// }
//}
if (bindex.x < 0)
{ {
// Pick up the mirroring value. // Pick up the mirroring value.
if ((i_dst < bsize) && ((bindex.x == 3) || (bindex.x ==1))) if ((i_dst < bsize))
{ {
boundloc = bsize; //Location of central border point. boundloc = bsize; //Location of central border point.
i_src = 2*boundloc - i_dst; i_src = 2*boundloc - i_dst;
} else if ((i_dst >= DCONST(AC_nx_min) - bsize) && ((bindex.x == 2) || (bindex.x ==1))) } else if ((i_dst >= DCONST(AC_nx_min) - bsize))
{ {
boundloc = DCONST(AC_nx_min) - bsize - 1; //Location of central border point. boundloc = DCONST(AC_nx_min) - bsize - 1; //Location of central border point.
i_src = 2*boundloc - i_dst; i_src = 2*boundloc - i_dst;
} }
}
if (bindex.y != 0)
{
// Pick up the mirroring value. // Pick up the mirroring value.
if ((j_dst < bsize) && ((bindex.y == 3) || (bindex.y ==1))) if ((j_dst < bsize))
{ {
boundloc = bsize; //Location of central border point. boundloc = bsize; //Location of central border point.
j_src = 2*boundloc - j_dst; j_src = 2*boundloc - j_dst;
} else if ((j_dst >= DCONST(AC_nx_min) - bsize) && ((bindex.y == 2) || (bindex.y ==1))) } else if ((j_dst >= DCONST(AC_nx_min) - bsize))
{ {
boundloc = DCONST(AC_ny_min) - bsize - 1; //Location of central border point. boundloc = DCONST(AC_ny_min) - bsize - 1; //Location of central border point.
i_src = 2*boundloc - j_dst; i_src = 2*boundloc - j_dst;
} }
}
if (bindex.z != 0)
{
// Pick up the mirroring value. // Pick up the mirroring value.
if ((k_dst < bsize) && ((bindex.z == 3) || (bindex.z ==1))) if ((k_dst < bsize))
{ {
boundloc = bsize; //Location of central border point. boundloc = bsize; //Location of central border point.
k_src = 2*boundloc - k_dst; k_src = 2*boundloc - k_dst;
} else if ((i_dst >= DCONST(AC_nz_min) - bsize) && ((bindex.z == 2) || (bindex.z ==1))) } else if ((i_dst >= DCONST(AC_nz_min) - bsize))
{ {
boundloc = DCONST(AC_nz_min) - bsize - 1; //Location of central border point. boundloc = DCONST(AC_nz_min) - bsize - 1; //Location of central border point.
k_src = 2*boundloc - k_dst; k_src = 2*boundloc - k_dst;
} }
} }
const int src_idx = DEVICE_VTXBUF_IDX(i_src, j_src, k_src); 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); const int dst_idx = DEVICE_VTXBUF_IDX(i_dst, j_dst, k_dst);
vtxbuf[dst_idx] = sign*vtxbuf[src_idx]; // sign = 1 symmetric, sign = -1 antisymmetric vtxbuf[dst_idx] = sign*vtxbuf[src_idx]; // sign = 1 symmetric, sign = -1 antisymmetric
@@ -141,8 +177,10 @@ acKernelGeneralBoundconds(const cudaStream_t stream, const int3 start, const int
(unsigned int)ceil((end.y - start.y) / (float)tpb.y), (unsigned int)ceil((end.y - start.y) / (float)tpb.y),
(unsigned int)ceil((end.z - start.z) / (float)tpb.z)); (unsigned int)ceil((end.z - start.z) / (float)tpb.z));
int3 bc_top = config.int3_params[AC_bc_type_top]; int3 bc_top = {config.int_params[AC_bc_type_top_x], config.int_params[AC_bc_type_top_y],
int3 bc_bot = config.int3_params[AC_bc_type_bot]; config.int_params[AC_bc_type_top_z]};
int3 bc_bot = {config.int_params[AC_bc_type_bot_x], config.int_params[AC_bc_type_bot_y],
config.int_params[AC_bc_type_bot_z]};
if (bc_top.x == AC_BOUNDCOND_SYMMETRIC) if (bc_top.x == AC_BOUNDCOND_SYMMETRIC)
{ {
@@ -154,6 +192,19 @@ acKernelGeneralBoundconds(const cudaStream_t stream, const int3 start, const int
kernel_symmetric_boundconds<<<bpg, tpb, 0, stream>>>(start, end, vtxbuf, bindex, -1); kernel_symmetric_boundconds<<<bpg, tpb, 0, stream>>>(start, end, vtxbuf, bindex, -1);
ERRCHK_CUDA_KERNEL(); ERRCHK_CUDA_KERNEL();
} }
else if (bc_bot.x == AC_BOUNDCOND_PERIODIC)
{
kernel_periodic_boundconds<<<bpg, tpb, 0, stream>>>(start, end, vtxbuf);
ERRCHK_CUDA_KERNEL();
}
else
{
printf("ERROR: Boundary condition not recognized!\n");
printf("ERROR: bc_top = %i, %i, %i \n", bc_top.x, bc_top.y, bc_top.z);
printf("ERROR: bc_bot = %i, %i, %i \n", bc_bot.x, bc_bot.y, bc_bot.z);
return AC_FAILURE;
}
return AC_SUCCESS; return AC_SUCCESS;
} }