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

Dummy at the moment, but now the boundary condition kernel caller can see what vertex buffer name is in use.

Browse Source
This commit is contained in:
Miikka Vaisala
2020-11-23 15:43:52 +08:00
parent 2f0f6ceac2
commit 543c565e5d
3 changed files with 36 additions and 24 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
src/core/device.cc
View File

@@ -440,7 +440,7 @@ acDeviceGeneralBoundcondStep(const Device device, const Stream stream,
{ {
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], config, bindex); device->vba.in[vtxbuf_handle], vtxbuf_handle, config, bindex);
} }
AcResult AcResult

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

@@ -191,7 +191,8 @@ 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 AcMeshInfo config, const int3 bindex) AcReal* vtxbuf, const VertexBufferHandle vtxbuf_handle,
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),
@@ -203,28 +204,38 @@ acKernelGeneralBoundconds(const cudaStream_t stream, const int3 start, const int
int3 bc_bot = {config.int_params[AC_bc_type_bot_x], config.int_params[AC_bc_type_bot_y], 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]}; config.int_params[AC_bc_type_bot_z]};
if (bc_top.x == AC_BOUNDCOND_SYMMETRIC) //#if AC_MPI_ENABLED
{ // printf( "WARNING : NON-PERIODIC BOUNDARY CONDITIONS NOT SUPPORTER BY MPI! Only working at node level.\n");
kernel_symmetric_boundconds<<<bpg, tpb, 0, stream>>>(start, end, vtxbuf, bindex, 1); // return AC_FAILURE;
ERRCHK_CUDA_KERNEL(); //#endif
}
else if (bc_top.x == AC_BOUNDCOND_ANTISYMMETRIC) if ( vtxbuf_handle != -1) // This is a dummy to make swithing boundary condition with respect to more possible later
{ {
kernel_symmetric_boundconds<<<bpg, tpb, 0, stream>>>(start, end, vtxbuf, bindex, -1);
ERRCHK_CUDA_KERNEL(); if (bc_top.x == AC_BOUNDCOND_SYMMETRIC)
} {
else if (bc_top.x == AC_BOUNDCOND_PERIODIC) kernel_symmetric_boundconds<<<bpg, tpb, 0, stream>>>(start, end, vtxbuf, bindex, 1);
{ ERRCHK_CUDA_KERNEL();
kernel_periodic_boundconds<<<bpg, tpb, 0, stream>>>(start, end, vtxbuf); }
ERRCHK_CUDA_KERNEL(); else if (bc_top.x == AC_BOUNDCOND_ANTISYMMETRIC)
} {
else kernel_symmetric_boundconds<<<bpg, tpb, 0, stream>>>(start, end, vtxbuf, bindex, -1);
{ ERRCHK_CUDA_KERNEL();
printf("ERROR: Boundary condition not recognized!\n"); }
printf("ERROR: bc_top = %i, %i, %i \n", bc_top.x, bc_top.y, bc_top.z); else if (bc_top.x == AC_BOUNDCOND_PERIODIC)
printf("ERROR: bc_bot = %i, %i, %i \n", bc_bot.x, bc_bot.y, bc_bot.z); {
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_FAILURE;
} }
return AC_SUCCESS; return AC_SUCCESS;

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 AcMeshInfo config, const int3 bindex); AcReal* vtxbuf, const VertexBufferHandle vtxbuf_handle,
const AcMeshInfo config, const int3 bindex);
/** */ /** */