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

Compiles without the API funtion call.

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

9
include/astaroth.h
View File

@@ -49,6 +49,11 @@ typedef struct {
typedef enum { AC_SUCCESS = 0, AC_FAILURE = 1 } AcResult; 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, #define AC_GEN_ID(X) X,
typedef enum { typedef enum {
AC_FOR_RTYPES(AC_GEN_ID) // 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, AcResult acDeviceGeneralBoundcondStep(const Device device, const Stream stream,
const VertexBufferHandle vtxbuf_handle, const int3 start, 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, 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);
/** */ /** */

11
src/core/device.cc
View File

@@ -436,19 +436,19 @@ acDevicePeriodicBoundconds(const Device device, const Stream stream, const int3
AcResult AcResult
acDeviceGeneralBoundcondStep(const Device device, const Stream stream, acDeviceGeneralBoundcondStep(const Device device, const Stream stream,
const VertexBufferHandle vtxbuf_handle, const int3 start, 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); cudaSetDevice(device->id);
return acKernelGeneralBoundconds(device->streams[stream], start, end, return acKernelGeneralBoundconds(device->streams[stream], start, end,
device->vba.in[vtxbuf_handle], bindex); device->vba.in[vtxbuf_handle], config, bindex);
} }
AcResult AcResult
acDeviceGeneralBoundconds(const Device device, const Stream stream, const int3 start, 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) { 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; return AC_SUCCESS;
} }
@@ -1435,6 +1435,8 @@ acGridStoreMesh(const Stream stream, AcMesh* host_mesh)
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
acGridGeneralBoundconds(const Device device, const Stream stream) acGridGeneralBoundconds(const Device device, const Stream stream)
{ {
@@ -1475,6 +1477,7 @@ acGridGeneralBoundconds(const Device device, const Stream stream)
return AC_SUCCESS; return AC_SUCCESS;
} }
*/
/* /*
// Unused // Unused

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

@@ -1,9 +1,5 @@
#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)
{ {
@@ -138,19 +134,22 @@ 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 int3 bindex) AcReal* vtxbuf, const AcMeshInfo config, 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_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<<<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_bc_type) == BOUNDCOND_ASYM) else if (bc_bot.x == AC_BOUNDCOND_ANTISYMMETRIC)
{ {
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();

3
src/core/kernels/kernels.h
View File

@@ -45,7 +45,8 @@ AcResult acKernelPeriodicBoundconds(const cudaStream_t stream, const int3 start,
AcReal* vtxbuf); AcReal* vtxbuf);
/** */ /** */
AcResult acKernelGeneralBoundconds(const cudaStream_t stream, const int3 start, const int3 end, 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); AcResult acKernelDummy(void);