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

Figuring out compilations.

Browse Source
This commit is contained in:
Miikka Vaisala
2020-11-20 11:58:15 +08:00
parent d060e67bec
commit cb15668f2d
4 changed files with 21 additions and 4 deletions
Download Patch File Download Diff File Expand all files Collapse all files

5
acc/mhd_solver/stencil_kernel.ac
View File

@@ -20,9 +20,12 @@
uniform int AC_max_steps; uniform int AC_max_steps;
uniform int AC_save_steps; uniform int AC_save_steps;
uniform int AC_bin_steps; uniform int AC_bin_steps;
uniform int AC_bc_type;
uniform int AC_start_step; uniform int AC_start_step;
// In3 params
uniform int3 AC_bc_type_bot;
uniform int3 AC_bc_type_top;
// Real params // Real params
uniform Scalar AC_dt; uniform Scalar AC_dt;
uniform Scalar AC_max_time; uniform Scalar AC_max_time;

6
include/astaroth.h
View File

@@ -304,7 +304,13 @@ AcResult acGridStoreMesh(const Stream stream, AcMesh* host_mesh);
AcResult acGridIntegrate(const Stream stream, const AcReal dt); AcResult acGridIntegrate(const Stream stream, const AcReal dt);
/** */ /** */
/* MV: Commented out for a while, but save for the future when standalone_MPI
works with periodic boundary conditions.
AcResult
acGridIntegrateNonperiodic(const Stream stream, const AcReal dt)
AcResult acGridIntegrateNonperiodic(const Stream stream, const AcReal dt); AcResult acGridIntegrateNonperiodic(const Stream stream, const AcReal dt);
*/
/** */ /** */
AcResult acGridPeriodicBoundconds(const Stream stream); AcResult acGridPeriodicBoundconds(const Stream stream);

4
src/core/device.cc
View File

@@ -1907,6 +1907,8 @@ acGridIntegrate(const Stream stream, const AcReal dt)
return AC_SUCCESS; return AC_SUCCESS;
} }
/* MV: Commented out for a while, but save for the future when standalone_MPI
works with periodic boundary conditions.
AcResult AcResult
acGridIntegrateNonperiodic(const Stream stream, const AcReal dt) acGridIntegrateNonperiodic(const Stream stream, const AcReal dt)
{ {
@@ -2130,6 +2132,8 @@ acGridIntegrateNonperiodic(const Stream stream, const AcReal dt)
return AC_SUCCESS; return AC_SUCCESS;
} }
*/
AcResult AcResult
acGridPeriodicBoundconds(const Stream stream) acGridPeriodicBoundconds(const Stream stream)

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

@@ -1,5 +1,9 @@
#pragma once #pragma once
// Temporary defines untils we can figure out smarter swithches.
#define BOUNDCOND_SYM 1
#define BOUNDCOND_ASYM 2
static __global__ void static __global__ void
kernel_symmetric_boundconds(const int3 start, const int3 end, AcReal* vtxbuf, const int3 bindex, const int sign) kernel_symmetric_boundconds(const int3 start, const int3 end, AcReal* vtxbuf, const int3 bindex, const int sign)
{ {
@@ -134,19 +138,19 @@ acKernelPeriodicBoundconds(const cudaStream_t stream, const int3 start, const in
AcResult AcResult
acKernelGeneralBoundconds(const cudaStream_t stream, const int3 start, const int3 end, acKernelGeneralBoundconds(const cudaStream_t stream, const int3 start, const int3 end,
AcReal* vtxbuf, const int bindex); AcReal* vtxbuf, const int3 bindex)
{ {
const dim3 tpb(8, 2, 8); const dim3 tpb(8, 2, 8);
const dim3 bpg((unsigned int)ceil((end.x - start.x) / (float)tpb.x), const dim3 bpg((unsigned int)ceil((end.x - start.x) / (float)tpb.x),
(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));
if (DCONST(AC_bype) == BOUNDCOND_SYM) if (DCONST(AC_bc_type) == BOUNDCOND_SYM)
{ {
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 (DCONST(AC_bype) == BOUNDCOND_ASYM) else if (DCONST(AC_bc_type) == BOUNDCOND_ASYM)
{ {
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();