diff --git a/include/astaroth.h b/include/astaroth.h index ace32c0..76f5174 100644 --- a/include/astaroth.h +++ b/include/astaroth.h @@ -49,6 +49,11 @@ typedef struct { typedef enum { AC_SUCCESS = 0, AC_FAILURE = 1 } AcResult; +// Neming the associated number of the boundary condition types +typedef enum { AC_BOUNDCOND_PERIODIC = 0, + AC_BOUNDCOND_SYMMETRIC = 1, + AC_BOUNDCOND_ANTISYMMETRIC = 2 } AcBoundcond; + #define AC_GEN_ID(X) X, typedef enum { AC_FOR_RTYPES(AC_GEN_ID) // @@ -577,11 +582,11 @@ AcResult acDevicePeriodicBoundconds(const Device device, const Stream stream, co /** */ AcResult acDeviceGeneralBoundcondStep(const Device device, const Stream stream, const VertexBufferHandle vtxbuf_handle, const int3 start, - const int3 end, const int3 bindex); + const int3 end, const AcMeshInfo config, const int3 bindex); /** */ AcResult acDeviceGeneralBoundconds(const Device device, const Stream stream, const int3 start, - const int3 end, const int3 bindex); + const int3 end, const AcMeshInfo config, const int3 bindex); /** */ diff --git a/src/core/device.cc b/src/core/device.cc index ad46148..9130f34 100644 --- a/src/core/device.cc +++ b/src/core/device.cc @@ -436,19 +436,19 @@ acDevicePeriodicBoundconds(const Device device, const Stream stream, const int3 AcResult acDeviceGeneralBoundcondStep(const Device device, const Stream stream, const VertexBufferHandle vtxbuf_handle, const int3 start, - const int3 end, const int3 bindex) + const int3 end, const AcMeshInfo config, const int3 bindex) { cudaSetDevice(device->id); return acKernelGeneralBoundconds(device->streams[stream], start, end, - device->vba.in[vtxbuf_handle], bindex); + device->vba.in[vtxbuf_handle], config, bindex); } AcResult acDeviceGeneralBoundconds(const Device device, const Stream stream, const int3 start, - const int3 end, const int3 bindex) + const int3 end, const AcMeshInfo config, const int3 bindex) { for (int i = 0; i < NUM_VTXBUF_HANDLES; ++i) { - acDeviceGeneralBoundcondStep(device, stream, (VertexBufferHandle)i, start, end, bindex); + acDeviceGeneralBoundcondStep(device, stream, (VertexBufferHandle)i, start, end, config, bindex); } return AC_SUCCESS; } @@ -1435,6 +1435,8 @@ acGridStoreMesh(const Stream stream, AcMesh* host_mesh) return AC_SUCCESS; } +/* MV: Commented out for a while, but save for the future when standalone_MPI + works with periodic boundary conditions. AcResult acGridGeneralBoundconds(const Device device, const Stream stream) { @@ -1475,6 +1477,7 @@ acGridGeneralBoundconds(const Device device, const Stream stream) return AC_SUCCESS; } +*/ /* // Unused diff --git a/src/core/kernels/boundconds.cuh b/src/core/kernels/boundconds.cuh index 01e89fa..64b4a1e 100644 --- a/src/core/kernels/boundconds.cuh +++ b/src/core/kernels/boundconds.cuh @@ -1,9 +1,5 @@ #pragma once -// Temporary defines untils we can figure out smarter swithches. -#define BOUNDCOND_SYM 1 -#define BOUNDCOND_ASYM 2 - static __global__ void kernel_symmetric_boundconds(const int3 start, const int3 end, AcReal* vtxbuf, const int3 bindex, const int sign) { @@ -138,19 +134,22 @@ acKernelPeriodicBoundconds(const cudaStream_t stream, const int3 start, const in AcResult acKernelGeneralBoundconds(const cudaStream_t stream, const int3 start, const int3 end, - AcReal* vtxbuf, const int3 bindex) + AcReal* vtxbuf, const AcMeshInfo config, const int3 bindex) { const dim3 tpb(8, 2, 8); 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.z - start.z) / (float)tpb.z)); - if (DCONST(AC_bc_type) == BOUNDCOND_SYM) + int3 bc_top = config.int3_params[AC_bc_type_top]; + int3 bc_bot = config.int3_params[AC_bc_type_bot]; + + if (bc_top.x == AC_BOUNDCOND_SYMMETRIC) { kernel_symmetric_boundconds<<>>(start, end, vtxbuf, bindex, 1); ERRCHK_CUDA_KERNEL(); } - else if (DCONST(AC_bc_type) == BOUNDCOND_ASYM) + else if (bc_bot.x == AC_BOUNDCOND_ANTISYMMETRIC) { kernel_symmetric_boundconds<<>>(start, end, vtxbuf, bindex, -1); ERRCHK_CUDA_KERNEL(); diff --git a/src/core/kernels/kernels.h b/src/core/kernels/kernels.h index 980bafa..bc94185 100644 --- a/src/core/kernels/kernels.h +++ b/src/core/kernels/kernels.h @@ -45,7 +45,8 @@ AcResult acKernelPeriodicBoundconds(const cudaStream_t stream, const int3 start, AcReal* vtxbuf); /** */ AcResult acKernelGeneralBoundconds(const cudaStream_t stream, const int3 start, const int3 end, - AcReal* vtxbuf, const int bound_direction); + AcReal* vtxbuf, const AcMeshInfo config, const int3 bindex); + /** */ AcResult acKernelDummy(void);